Sau khi viết xong bài đầu tiên, mình tiếp tục tìm hiểu sâu hơn về ma trận và cảm thấy cần phải viết thêm một (vài) bài nữa vì có quá nhiều thứ hay ho khi đi sâu vào.
Bạn đang xem: Lập trình song song trên gpu
Trong phần này, mình giới thiệu về phương pháp sử dụng GPU để thực hiện phép tính nhân 2 ma trận (GPU Calculation), thay cho việc dùng CPU truyền thống, để tăng tốc độ tính toán.
Tại sao lại là GPU? có cả việc tính toán trên GPU sao?
Câu trả lời là có, thưa các bạn. Chúng ta đều biết GPU là một thành phần trong card đồ họa, thường được giới game thủ mua về cắm vô cho máy tính xử lý đồ họa nhanh hơn. Vậy thao tác xử lý đồ họa ở đây là gì?
Để hiểu được chức năng của GPU, hãy cùng lái chủ đề một tí sang lĩnh vực đồ họa máy tính, viết đến đây mình thấy hơi bị máu vì dù đã từ bỏ ngành game được hơn 1 năm rồi nhưng giờ vẫn có cơ hội quay lại viết tiếp về nó, hê hê, lạc đề tí thôi, các bạn cũng yên tâm vì mình sẽ không để bài viết này thành bài hướng dẫn làm game gì đâu =)))
Máy tính tạo ra hình ảnh như thế nào?
Một hình ảnh trên máy tính được cấu tạo thành từng pixel, điều này hẳn ai cũng biết, nhưng làm thế nào để máy tính hiểu và vẽ được một hình ảnh thành một tổ hợp các pixel đưa lên màn hình qua sự chỉ đạo của chương trình (application)?
Trong trường hợp là game thì các dữ liệu hình ảnh (3d mesh, texture,...) sẽ được load từ đĩa (hard disk, file,...) và truyền vào GPU. Tại đây, GPU sẽ tiến hành phân tích và thực hiện những công việc tính toán cần thiết để chuyển chúng thành tập hợp các pixel gửi trả về CPU để vẽ ra màn hình máy tính (IO devices).
Quá trình tính toán đó ở trong GPU, được gọi là Rendering Pipeline, trong hình minh họa ở trên thì dây chuyền vận chuyển các khối hộp pixel kia chính là pipeline.
Rendering Pipeline
Một vật thể 3D trong máy tính được cấu thành từ nhiều đỉnh, gọi là vertex, mảng các vertex được gọi là vertices array.
Hình trên cho thấy một khối hộp được tạo từ mảng $v$ gồm các đỉnh $v_{0}, v_{1}, \ldots, v_{7}$
Trong máy tính, mọi hình ảnh từ đơn giản hay phức tạp đều được tạo thành từ các vertices như vậy.
Render Pipeline là quá trình xảy ra trong GPU, khi nó nhận một tập hợp các vertices, trải qua các bước xử lý kết nối các vertices lại với nhau (vertex processing) thành các khối 3D cơ bản (thường là các khối tam giác, gọi là 3D primitives), sau đó phân tích (rasterization) các hình khối 3D này thành từng mảnh nhỏ (fragments), thực hiện việc tính toán tô màu, đổ bóng,... cho các khối này rồi ánh xạ nó thành mảng các pixel tương ứng với màn hình máy tính, gửi trả về CPU để đưa lên màn hình.
Vì quy trình này diễn ra liên tục với tốc độ cực kì cao để đảm bảo cho việc hiển thị hình ảnh lên màn hình diễn ra xuyên suốt và mượt mà, nên GPU phải làm việc với một khối lượng các phép tính cực kì lớn. Nói như vậy thì không có nghĩa là CPU không có khả năng dựng hình ảnh đồ họa trên máy tính, bản thân CPU cũng có thể làm được nhưng sẽ không đạt được performance như GPU, lý do tại sao xin mời các bạn xem tiếp.
GPU vs CPU
Hình sau so sánh kiến trúc của một CPU so với một GPU:
Các bạn có thể thấy khối tính toán của GPU nhiều hơn CPU rất nhiều lần, vì thế làm cho GPU có khả năng tính toán cao hơn CPU.
Thêm một hình ảnh nữa để so sánh sự lợi hại của GPU, nếu bạn có một máy tính có vi xử lý 4 nhân thì đây là sự khác biệt đối với GPU:
Nhưng nếu sử dụng GPU lợi hại như vậy thì tại sao người ta không dùng GPU để thay thế CPU luôn? Câu trả lời là rất khó, vì việc debug trên GPU không đơn giản như với CPU, thêm nữa là GPU code chạy song song (parallel) nên cần phải có các phương pháp để phân vùng dữ liệu, liên lạc, đồng bộ hóa giữa các core nên các thuật toán xử lý trên GPU cực kì phức tạp.
Đoạn video sau đây minh họa sự khác biệt về việc dùng CPU và GPU để vẽ hình ảnh, con robot màu vàng là CPU còn cái máy bự chà bá kia là sức mạnh thực sự của một GPU.
GPU chỉ dùng làm đồ họa?
Với sự phát triển của lĩnh vực công nghệ thông tin hiện nay, những công việc yêu cầu khả năng tính toán cực kì cao đã bắt đầu cần đến các hệ thống máy tính có khả năng xử lý mạnh hơn, và người ta đã bắt đầu tìm đến sự trợ giúp của các GPU trong công việc tính toán, không phải chỉ riêng CPU nữa.
Một ứng dụng thực tế có lẽ nhiều người biết đến nhất đó là việc sử dụng GPU để "đào" Bitcoin.
Một ứng dụng khác thường thấy nữa là trong ngành Machine Learning, người ta thường sử dụng GPU để gia tăng tốc độ tính toán đối với các vector dữ liệu, vì đây cũng là một ngành phải xử lý với lượng dữ liệu cực kì lớn.
Gói EC2 P2 của AWS cũng hỗ trợ GPU https://aws.amazon.com/blogs/aws/new-p2-instance-type-for-amazon-ec2-up-to-16-gpus/
Tính toán với GPU bằng cách nào?
Có rất nhiều cách để sử dụng GPU làm công cụ tính toán, bạn có thể khởi tạo một ứng dụng đồ họa, chuyển đổi vector dữ liệu cần tính toán thành các ma trận texture rồi nạp vào GPU, viết các Shader program (là một dạng chương trình chạy trên GPU, trong quá trình fragment processing của rendering pipeline) để tính toán và trả dữ liệu về lại CPU.
Tuy nhiên cách làm trên khá là phức tạp và dư thừa (phải khởi tạo môi trường đồ họa, không cần thiết, và đó là cách làm thủ công, không có gì đảm bảo, và bị giới hạn bởi khả năng của Shader program.
Ngày nay chúng ta có các giải pháp khác, có thể nói là "native" hơn, đó là Open
CL và CUDA.
CUDA
Khi mới được ra mắt, CUDA là tên viết tắt của Compute Unified Device Architecture tuy nhiên về sau thì Nvidia bỏ hẳn tên gọi này.
CUDA là một nền tảng lập trình song song (parallel computing platform), cung cấp cho chúng ta các API trên các ngôn ngữ như C/C++, Fortran, và về sau còn xuất hiện thêm các wrapper cho Java, Ruby, Python, Haskell... để giao tiếp với các GPU của Nvidia và thực hiện công việc tính toán giống như trên CPU.
Quy trình hoạt động của một ứng dụng CUDA khá là giống với rendering pipeline:
GPU thực hiện tính toán song song trên các core của nó
Tính xong GPU gửi trả dữ liệu về lại cho bộ nhớ chính
CUDA có một vài nhược điểm nhất định, tuy nhiên điểm yếu lớn nhất của nó là được làm ra bởi Nvidia nên chỉ hỗ trợ chạy trên phần cứng của Nvidia, và không ai muốn công nghệ mình sử dụng lại bị phụ thuộc vào một hãng sản xuất nhất định nào hết (trừ trường hợp Apple =))) vậy cho nên người ta bắt đầu tìm đến một giải pháp tổng quát hơn, đó là Open
CL
Open
CL
Open Computing Language viết tắt là Open
CL là một framework giúp các lập trình viên phát triển các ứng dụng mà chương trình có thể chạy được trên nhiều bộ xử lý như CPU, GPU, FPGA,...
Open
CL là một chuẩn mở được Apple tạo ra và trao quyền phát triển cho tổ chức Khronos Group trứ danh (tổ chức đang nắm quyền phát triển Open
GL). Open
CL được rất nhiều hãng phần cứng hỗ trợ như là AMD, Apple, ARM, IBM, Imagination Technologies, Intel, Nvidia, Qualcomm, Samsung,... điều này khiến cho Open
CL trở thành một tiêu chuẩn chung cho cả ngành công nghiệp, giống như Open
GL vậy.
Và trong bài viết này chúng ta sẽ sử dụng Open
CL để implement thuật toán nhân ma trận trên GPU. Để có thể implement được thì đầu tiên phải tìm hiểu rõ hơn về Open
CL và cách lập trình sử dụng Open
CL.
Lập trình GPU với Open
CL
Đầu tiên chúng ta cần hiểu các khái niệm sử dụng trong Open
CL.
Các thuật ngữ trong Open
CL
Các thành phần tham gia việc tính toán trong một ứng dụng Open
CL là các Compute Device, gọi là các thiết bị tính toán. Và các thiết bị tính toán này có thể là các CPU, GPU hoặc các thiết bị khác có tính năng tương tự bên trong máy tính.
Trong mỗi một Compute Device chúng ta có nhiều Compute Unit, là các core bên trong nó.
Và trong mỗi Compute Unit (core) ta có nhiều Processing Element trực tiếp nhận lệnh để tính toán.
Mô hình làm việc trong Open
CL
Mô hình xử lý của Open
CL gồm có một chương trình Host có nhiệm vụ điều khiển và liên lạc với nhiều Compute Device theo sơ đồ sau:
Khi thực hiện xử lý, Host sẽ khởi tạo trên mỗi Compute Device một Context chứa các thành phần như là:
Thông tin device đang chạyKernel (là đơn vị nhỏ nhất của quá trình xử lý, như là một hàm)Program Object (là đoạn code implement của Kernel)Command Queues (là queue các lệnh được truyền vào device từ host)Memory Object (là đối tượng chứa dữ liệu sẽ được truyền/nhận qua lại giữa device và host)
Thao tác copy dữ liệu từ Host tới Compute Device gọi là thao tác ghi (writing), và thao tác copy dữ liệu ngược lại từ Compute Device về Host gọi là đọc (reading).
Đây là flow của một chương trình Open
CL thường thấy:
Có thể thấy mô hình trên có phần phức tạp hơn CUDA. Nhưng làm việc với Open
CL linh hoạt hơn và không bị giới hạn về mặt thiết bị như CUDA.
Việc khởi động một kernel rất tốn kém, vì thế lời khuyên được đưa ra là nên thiết kế mỗi kernel làm nhiều việc nhất có thể.
Mô hình quản lý bộ nhớ của Open
CL
Implement một chương trình Open
CL đơn giản
Vậy coi như xong phần lý thuyết, giờ chúng ta đi vào phần thực hành cơ bản để hiểu kĩ hơn về cách lập trình trên GPU.
Đề bài sẽ là: Implement chương trình tính tổng từ 0 đến 100 triệu
Implement trên CPU
Với cách này thì không có gì để bàn rồi, đơn giản chúng ta sẽ implement một vòng for rồi cho nó chạy từ 0 đến 100 triệu, cứ thế mà cộng số vào.
for (unsigned long long i = 0; i Chúng ta sẽ thêm vào một đoạn code để đo thời gian chương trình bỏ ra để thực hiện phép tính trên.
clock_t begin = clock();// Do somethingclock_t end = clock();double runtime = (double)(end - begin) / CLOCKS_PER_SEC;Hàm clock() lấy thời gian hiện tại, có trong thư viện time.h, mục đích của đoạn code trên là lấy thời gian ở vị trí bắt đầu xử lý, và thời gian ở vị trí sau khi kết thúc xử lý, trừ 2 số đó với nhau ta có được thời gian cần để chạy chương trình.
Đoạn implement đầy đủ sẽ như sau:
#include #include typedef unsigned long long fuckin_large;int main() { clock_t begin = clock(); fuckin_large val = 0; for (fuckin_large j = 0; j Chạy thử chương trình trên, output sẽ có dạng như sau:
Result: 4999999950000000Runtime: 0.225600ms
Dòng đầu là con số kết quả sau khi đã tính ra, dòng thứ 2 cho biết thời gian cần để thực hiện chương trình.
Mình có chạy thử chương trình này trên các máy Mac
Book Pro từ 2012 tới 2015, trung bình mất tầm 0.2 mili giây để tính ra kết quả.
Tiếp theo hãy thử implement chương trình này với cách xử lý tương tự nhưng chạy trên GPU xem tốc độ được cải thiện như thế nào nhé.
Implement trên GPU
Để implement một chương trình dùng Open
CL thì ta cần implement 2 phần: Host program và Kernel program.
Host là một chương trình C/C++ có nhiệm vụ đọc file chương trình Kernel (đuôi *.cl) và làm các công việc khác như là khởi tạo context, chọn device, build program object, memory object, thực thi kernel,... như đã nói ở phần trên.
Host ProgramBước 1: Khởi tạo
Đầu tiên Host cần khởi tạo các biến để chứa các đối tượng cần thiết như là Device, Context, Command Queue,...:
cl_device_id device_id = NULL;cl_context context = NULL;cl_command_queue command_queue = NULL;cl_mem memobj = NULL;cl_program program = NULL;cl_kernel kernel = NULL;cl_platform_id platform_id = NULL;cl_uint ret_num_devices;cl_uint ret_num_platforms;cl_int ret;Bước 2: Đọc Kernel Source
Tiếp theo chúng ta đọc file Kernel source và lưu vào một file pointer:
FILE *fp;char file
Name<> = "./sum.cl";char *source_str;size_t source_size;fp = fopen(file
Name, "r");if (!fp) { fprintf(stderr, "Failed to load kernel\n"); exit(1);}source_str = (char*)malloc(MAX_SOURCE_SIZE);source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);fclose(fp);Trong ví dụ này thì file Kernel của chúng ta là sum.cl
Bước 3: Khởi tạo Context và các thành phần liên quan
Sau khi đã có source code của Kernel, ta tiến hành chọn device và khởi tạo context, biên dịch source code của kernel thành Program Object
ret = cl
Get
Platform
IDs(1, &platform_id, &ret_num_platforms);ret = cl
Get
Device
IDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);context = cl
Create
Context(NULL, 1, &device_id, NULL, NULL, &ret);command_queue = cl
Create
Command
Queue(context, device_id, 0, &ret);program = cl
Create
Program
With
Source(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);ret = cl
Build
Program(program, 1, &device_id, NULL, NULL, NULL);kernel = cl
Create
Kernel(program, "hello", &ret);Bước 4: Cấp phát bộ nhớ trên Compute Device
Sau đó cấp phát bộ nhớ cho Compute Device trong Context thông qua hàm cl
Create
Buffer và cl
Set
Kernel
Arg để chỉ định vùng nhớ này tương ứng với tham số nào khi truyền vào kernel.
memobj = cl
Create
Buffer(context, CL_MEM_READ_WRITE, sizeof(cl_mem), NULL, &ret);ret = cl
Set
Kernel
Arg(kernel, 0, sizeof(cl_mem), (void*)&memobj);Ngoài ra, chúng ta có thể truyền giá trị vào cho các tham số thông qua hàm cl
Enqueue
Write
Buffer:
ret = cl
Enqueue
Write
Buffer(command_queue, memobj_A, CL_TRUE, 0, MATRIX_SIZE * MATRIX_SIZE * sizeof(cl_mem), A, 0, NULL, NULL);Bước 5: Chạy Kernel và Đọc kết quả
Đến bước này chúng ta sử dụng hàm cl
Enqueue
Task để gán Command Queue đã tạo vào cho Kernel. Sau khi được gán thì Kernel sẽ khởi động và bắt đầu xử lý trên các Processing Elements trên Compute Device (ở đây có thể là GPU của chúng ta).
ret = cl
Enqueue
Task(command_queue, kernel, 0, NULL, NULL);Để lấy dữ liệu ra sau khi Kernel hoàn thành việc tính toán, chúng ta dùng hàm cl
Enqueue
Read
Buffer:
ret = cl
Enqueue
Read
Buffer(command_queue, memobj, CL_TRUE, 0, sizeof(cl_mem), val, 0, NULL, NULL);Giá trị CL_TRUE truyền vào hàm này là cho tham số cl_bool blocking_read, tham số này cho biết hàm cl
Enqueue
Read
Buffer sẽ block không cho Host program thực hiện tiếp cho đến khi nó nhận được giá trị trả về sau khi Kernel hoàn thành.
Đây cũng là bước mà chúng ta cần sử dụng hàm clock() để tính toán thời gian xử lý của Kernel.
clock_t begin = clock();// Do somethingclock_t end = clock();double runtime = (double)(end - begin) / CLOCKS_PER_SEC;Bước 6: Giải phóng bộ nhớ
Kết thúc chương trình, chúng ta giải phóng tất cả các đối tượng đã tạo ra để tránh bị memory leak.
ret = cl
Flush(command_queue);ret = cl
Finish(command_queue);ret = cl
Release
Kernel(kernel);ret = cl
Release
Program(program);ret = cl
Release
Mem
Object(memobj);ret = cl
Release
Command
Queue(command_queue);ret = cl
Release
Context(context);free(source_str);Đây là đoạn code implement đầy đủ của Host program làm công việc tính tổng theo đề bài đã cho:
#include #include #include #include #define MAX_SOURCE_SIZE (0x100000)int main() { cl_device_id device_id = NULL; cl_context context = NULL; cl_command_queue command_queue = NULL; cl_mem memobj = NULL; cl_program program = NULL; cl_kernel kernel = NULL; cl_platform_id platform_id = NULL; cl_uint ret_num_devices; cl_uint ret_num_platforms; cl_int ret; cl_ulong val<1>; FILE *fp; char file
Name<> = "./sum.cl"; char *source_str; size_t source_size; fp = fopen(file
Name, "r"); if (!fp) { fprintf(stderr, "Failed to load kernel\n"); exit(1); } source_str = (char*)malloc(MAX_SOURCE_SIZE); source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp); fclose(fp); ret = cl
Get
Platform
IDs(1, &platform_id, &ret_num_platforms); ret = cl
Get
Device
IDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices); context = cl
Create
Context(NULL, 1, &device_id, NULL, NULL, &ret); command_queue = cl
Create
Command
Queue(context, device_id, 0, &ret); memobj = cl
Create
Buffer(context, CL_MEM_READ_WRITE, sizeof(cl_mem), NULL, &ret); program = cl
Create
Program
With
Source(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret); ret = cl
Build
Program(program, 1, &device_id, NULL, NULL, NULL); kernel = cl
Create
Kernel(program, "hello", &ret); ret = cl
Set
Kernel
Arg(kernel, 0, sizeof(cl_mem), (void*)&memobj); clock_t begin = clock(); ret = cl
Enqueue
Task(command_queue, kernel, 0, NULL, NULL); ret = cl
Enqueue
Read
Buffer(command_queue, memobj, CL_TRUE, 0, sizeof(cl_mem), val, 0, NULL, NULL); clock_t end = clock(); double runtime = (double)(end - begin) / CLOCKS_PER_SEC; ret = cl
Flush(command_queue); ret = cl
Finish(command_queue); ret = cl
Release
Kernel(kernel); ret = cl
Release
Program(program); ret = cl
Release
Mem
Object(memobj); ret = cl
Release
Command
Queue(command_queue); ret = cl
Release
Context(context); printf("Result: %llu\n", val<0>); printf("Runtime: %lfms\n", runtime); free(source_str); return 0;}Kernel Program
Tiếp đến chúng ta implement Kernel, đây là đoạn chương trình sẽ được Host nạp vào Program Object để thực hiện xử lý trên GPU.
Chương trình này sẽ được định nghĩa bằng các từ khóa kernel. Ta truyền vào một biến con trỏ để tiếp nhận dữ liệu đầu ra, chương trình này sử dụng phương pháp tính bằng cách chạy một vòng lặp for, có thể implement như sau:
kernel void hello(global ulong *val) { size_t i = get_global_id(0); for (ulong j = 0; j Biến i trong chương trình trên là id của tiến trình kernel được khởi tạo trong lúc thực thi, chúng ta cần phải lấy chính xác id này để gửi trả dữ liệu về cho Host.
Đặc điểm của kernel là các biến, tham số truyền vào đều mutable và chúng ta không cần return dữ liệu (vì ở host đã có một bước đọc dữ liệu từ memory object ra rồi).
Hy vọng qua bài viết này, các bạn đã nắm được phần nào khái niệm về việc sử dụng GPU để cải thiện hiệu quả tính toán, và lý do tại sao phải dùng GPU.
Trong phần tiếp theo, mình sẽ giới thiệu cách implement thuật toán nhân ma trận trên GPU.
Xin cảm ơn các bạn đã theo dõi bài viết, và hẹn gặp lại các bạn trong các bài viết sắp tới.
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.A tag already exists with the provided branch name. Many Git commands accept both tag and branch names, so creating this branch may cause unexpected behavior. Are you sure you want to create this branch?
Launching Visual Studio Code Your codespace will open once ready.
There was a problem preparing your codespace, please try again.
Xem thêm: It - Chú Hề Ma Quái
Lập trình song song GPU: Đồ án cuối kỳ 1. Thông tin nhóm 2. Yêu cầu của đồ án 3. Kế hoạch thực hiện (dự kiến) 4. Họp nhóm và phân công công việc 5. Kết quả các baseline + Hình ảnh chụp Baseline 1 (Version1.cu) Baseline 2 (Version1.cu) Baseline 3 (Version2.cu) Baseline 4 (Version3.cu) 6. Quy trình tối ưu thuật toán Radix Sort Bảng đo thời gian 6.1. Lần tối ưu hóa 1 Phân tích Thiết kế Cài đặt 6.2. Lần tối ưu hóa 2 Phân tích Thiết kế Cài đặt 6.3 Lần tối ưu hóa 3 Phân tích Thiết kế Cài đặt 6.4 Lần tối ưu hóa 4 Phân tích Thiết kế Cài đặt 6.5 Lần tối ưu hóa 5 Phân tích Thiết kế Cài đặt 6.6 Lần tối ưu hóa 6 Phân tích Thiết kế Cài đặt 6.7 Lần tối ưu hóa 7 Phân tích Thiết kế Cài đặt 6.8 Lần tối ưu hóa 8 Phân tích Thiết kế Cài đặt
README.md
Lập trình song song GPU: Đồ án cuối kỳ
Lập trình song song GPU: Đồ án thuật toán sắp xếp Radix Sort
1. Thông tin nhóm
| 1 | Trần Nhật Huy | 1612272 |
| 2 | Nguyễn Thành Đạt | 1612088 |
2. Yêu cầu của đồ án
Đồ án sẽ gồm có 4 phiên bản baseline:Cài đặt tuần tự thuật toán Radix Sort tuần tự (đã làm ở BT4)Cài đặt song song 2 bước histogram và scan của thuật toán Radix Sort tuần tự (đã làm ở BT4)Cài đặt song song thuật toán Radix Sort với k = 1 bitCài đặt tuần tự ý tưởng của thuật toán Radix Sort song song trong bài báo của NVIDIA (đây là phiên bản mà lúc sau ta sẽ tập trung song song hóa và tối ưu hóa).
Đích đến: Tối ưu gần bằng thuật toán sắp xếp có sẵn trong thư viện Thrust
3. Kế hoạch thực hiện (dự kiến)
4. Họp nhóm và phân công công việcNội dung: Bàn về kế hoạch làm việc nhóm và tạo bảng kế hoạch thực hiện
Nội dung: Bàn về kế hoạch tối ưu hóa, lúc này nhóm vừa hoàn thành thuật toán nhưng thời gian thực thi của hàm scatter còn lâu
| Viết thuật toán Radix Sort tuần tự | Trần Nhật Huy |
| Viết thuật toán Radix Sort song song | Nguyễn Thành Đạt |
| Viết thuật toán Radix Sort song song với k = 1 | Trần Nhật Huy |
| Tối ưu thuật toán Radix Sort | Trần Nhật Huy và Nguyễn Thành Đạt |
| Làm slide | Nguyễn Thành Đạt |
| File làm việc nhóm | Trần Nhật Huy |
5. Kết quả các baseline + Hình ảnh chụp
Baseline 1 (Version1.cu)Baseline 2 (Version1.cu)
Baseline 3 (Version2.cu)
Baseline 4 (Version3.cu)
6. Quy trình tối ưu thuật toán Radix SortViệc tối ưu thuật toán sẽ chia làm 3 bước: Phân tích, Thiết kế và Cài đặt
Vì thầy sẽ đánh giá điểm dựa vào quy trình tối ưu thuật toán là có hợp lý hay không hơn là đánh giá vào thời gian chạy của thuật toán nên tất cả các bước tối ưu hóa sau khi cài đặt được thuật toán radix sort song song sẽ được ghi tiếp ở file này theo các bước trên
Việc làm như vậy sẽ dễ dàng trong việc viết báo cáo về sau
Bảng đo thời gian
| Baseline 1 | 852.525 | |||||
| Baseline 2 | 561.608 | |||||
| Baseline 3 | 1734.354 | |||||
| Baseline 4 (Tuần tự) | 19602.135 | |||||
| Parallel v1 | 13.332 | 0.782 | 0.179 | 0.529 | 117.443 | 171.905 |
| Parallel v2 | 13.324 | 0.780 | 0.179 | 0.531 | 73.052 | 128.528 |
| Parallel v3 | 13.504 | 0.738 | 0.174 | 0.514 | 70.114 | 123.828 |
| Parallel v4 | 13.327 | 0.762 | 0.175 | 0.529 | 59.814 | 113.641 |
| Parallel v5 | 13.318 | 0.769 | 0.181 | 0.529 | 60.569 | 114.792 |
| Parallel v6 | 13.324 | 0.761 | 0.179 | 0.531 | 60.1 | 113.534 |
| Parallel v7 | 4.457 | 0.777 | 0.18 | 0.53 | 59.952 | 104.299 |
| Thrust | 44.17 |
Ta sẽ tiến hành song song hóa các bước ở baseline 4 để giảm tổng thời gian thực hiện
Thiết kếThuật toán sẽ giống như được học trên lớp. Ở bước histogram
Kernel thì sau khi làm xong thì sẽ lưu xuống mảng 2 chiều với mỗi cột sẽ là một histogram cho dễ scan. Sau khi scan, ta sẽ chuyển vị lại cho dễ làm các bước sau. Ở bản này, bước scatter sẽ dùng một thread duy nhất trong block để chạy vì nhóm chưa tìm ra cách song song hóa.Cài đặt
Nhận xét: Tốc độ chạy nhanh hơn6.2. Lần tối ưu hóa 2Phân tích
Hiện tại, thuật toán chậm nhất ở phần Scatter. Lý do là vì nhóm đang sử atomic
Add để tính rank.Thiết kếNhóm sẽ viết lại thuật toán Scatter y như trong bài báo, sử dụng Radix Sort với k = 1 để tính rank của các phần tử trong Block
Cài đặt
Nhận xét: Tốc độ tăng lên đáng kể6.3 Lần tối ưu hóa 3Phân tích
Ở bước này, nhóm sẽ gán block
Size sao cho tận dụng được tối đa sức mạnh của GPU. Thực tế, bước này sẽ cần làm trước lần tối ưu hóa số 2 nhưng vì nhóm muốn hoàn thiện thuật toán giống như trong bài báo
Thiết kếNhóm sẽ gán block
Size = Số-thread-tối-đa-trên-một-SM / Số-lượng-block-tối-đa-trên-một-SMCài đặt
Nhận xét: Tốc độ chạy tăng lên6.4 Lần tối ưu hóa 4Phân tích
Ở bước này, nhóm sẽ vẫn tối ưu hóa Scatter kernel. Ở phiên bản Parallel_v3 thì trong Scatter Kernel nhóm sử dụng mảngsmem để lưu dữ liệu và kích thước của mảng này không hề nhỏ khi nhóm lưu tới 6 loại mảng.Song song với đó là nhóm sử dụng nhiều câu lệnh __syncthreads()Thiết kếThay vì smem trong Scatter kernel phải lưu 5 mảng khác nhau thì giờ đây nhóm sẽ thiếtkế lại:Phần 1 vẫn sẽ lưu input kích thước block
Dim.x phần tử
Phần 2 lưu một phần tử giả, gọi là dummy, có 1 phần tử
Phần 3 lưu chuỗi nhị phân lấy được, kích thước block
Dim.x phần tử
Phần 4 lưu chỉ số bắt đầu của mảng sau khi đã sắp xếp, kích thước 2 ^ n
Bits phần tử
Cài đặt
Nhận xét: Tốc độ chạy tăng lên6.5 Lần tối ưu hóa 5Phân tích
Ở lần này nhóm vẫn tiếp tục tối ưu hóa hàm Scatter kernel vì trong hàm này vẫn còn có nhiềucâu lệnh __syncthreads() dùng để debug và chưa load mảng scan
Histogram
Array
Tranpose vào smem
Thiết kếỞ lần trước thì smem gồm 4 phần dữ liệu, bây giờ ta sẽ thêm một phần chứa scan
Histogram
Array
Tranpose cho từng block, sẽ có 2 ^ n
Bits phần tử. Khi chép dữ liệu cần lưu ý rằng số thread có thể sẽ nhỏ hơn số n
Bins
Cài đặt
Nhận xét: Tốc độ chạy giảm đi một chút6.6 Lần tối ưu hóa 6Phân tích
Ở lần này, nhóm chưa suy nghĩ ra cách nào để tối ưu hàm Scatter kernel nữa. Do đó, nhóm sẽ thực hiện tối ưu ở điểm khác. Tổng thời gian chạy các hàm kernel chiếm phần lớn thời gian thực thi nhưng thời gian còn nằm ở những phần không thuộc các hàm kernel và nhóm sẽ tối ưu các hàm cấp phát bộ nhớ xem có giảm tổng thời gian thực thi hay không
Thiết kếỞ phần cấp phát bộ nhớ, hiện tại thì nhóm đang lãng phí tài nguyên khi có thể tái sử dụng d_hist
Arr thay vì tạo d_scan
Hist
Arr
Tranpose
Cài đặt
Nhận xét: Tốc độ chạy tăng lên6.7 Lần tối ưu hóa 7Phân tích
Với các phiên bản trước thì kernel tính hist sẽ sử dụng atomic
Add trên GMEM nên có tốc độ chậm
Thiết kếNhóm sẽ tiến hành tối ưu kernel tính hist bằng cách sử dụng atomic
Add trên SMEMCài đặt
Nhận xét: Tốc độ chạy tăng lên6.8 Lần tối ưu hóa 8Phân tích
Hàm Scatter Kernel hiện tại đang vừa sort với scatter trong một hàm. Do đó, hàm này đang sử dụng rất nhiều smem để lưu trữ mặc dù nhóm đã có tối ưu phần này.Thiết kếNhóm sẽ tách Scatter Kernel làm 2 hàm riêng biệt là sort
Block với scatter
Block. Phần code thì vẫn tương tự như Scatter Kernel nhưng nhóm muốn thử nghiệm xem nó có nhanh hơn hay không
Cài đặt
Nhận xét: ....