Programming, Network, Office, Hardware, Software, Design, Photography.. I am just a GEEK
Bài viết tiếp theo trong chuỗi bài viết về Lập trình song song trên GPU bằng CUDA.
Chúng ta đã tìm hiểu cách thực thi song song 2 bài toán cơ bản là Cộng 2 vector và Nhân 2 ma trận ở các bài viết trước
Trong bài này, húng ta sẽ cài đặt thuật toán Cộng các phần tử trong mảng hay có một tên gọi khác là Reduction trên CPU và GPU. Sau đó, chúng ta sẽ so sánh tốc độ thực thi trên một số cấu hình khác nhau.
Nội dung
Kích thước mảng: 2^13 + 1
Phát sinh ngẫu nhiên giá trị của các phần tử đầu vào trong [0, 1]
So sánh thời gian chạy với các kích thước block khác nhau
Thực hiện
Chúng ta cùng xem lại thuật toán tính tổng các phần tử trong mảng nào.
|
1 2 3 |
int sum = 0; for (int i = 0; i < n; i++) sum += array[i]; |
Vậy chúng ta sẽ song song hóa như thế nào?
Nhận thấy kết quả của tổng sẽ được cộng dồn sau mỗi bước nhảy, chúng ta thử chia đôi mảng ra và thực hiện cộng 2 phần lại xem sao
|
1 2 3 4 5 6 7 8 9 10 11 |
int reduce(int *data, int n) { if(n == 1) return data[0]; int stride = n/2; for(int i = 0; i < stride; i++) { data[i] += data[i + stride]; } return reduce(data, stride); } |
Ta áp dụng với CUDA tương tự bằng việc sử dụng thread để đảm nhiệm phần tử tương ứng.
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 |
#include <stdio.h> #include <stdlib.h> #include <math.h> #include <time.h> int reduceOnHost(int *data, int n) { if(n == 1) return data[0]; int stride = n/2; for(int i = 0; i < stride; i++) { data[i] += data[i + stride]; } return reduceOnHost(data, stride); } __global__ void reduceOnDevice(int *d_data, int *d_sum, int n) { int tid = threadIdx.x; int idx = blockIdx.x * blockDim.x + threadIdx.x; int *idata = d_data + blockIdx.x * blockDim.x; if(idx >= n) return; for(int stride = 1; stride < blockDim.x; stride *= 2) { if((tid % (2*stride)) == 0) { idata[tid] += idata[tid + stride]; } __syncthreads(); } if(tid == 0) d_sum[blockIdx.x] = idata[0]; } int main(int argc, char* argv[]) { int sizeArr[10] = {64, 256, 1024, 4069, 16384, 65536, 262144, 1048576, 4194304, 16777216}; for(int i = 0; i < 10; i++) { int n = sizeArr[i]; int *data, *data2, *sum; int *d_data, *d_sum; size_t bytes = n * sizeof(int); data = (int*)malloc(bytes); data2 = (int*)malloc(bytes); sum = (int*)malloc(bytes); cudaMalloc(&d_data, bytes); cudaMalloc(&d_sum, bytes); for(int j = 0; j < n; j++ ) { data[j] = 1; data2[j] = 1; } int hostSum = 0; clock_t begin, end; begin = clock(); hostSum = reduceOnHost(data, n); end = clock(); double duration = double(end-begin)/CLOCKS_PER_SEC * 1000; printf("(Vector size %d):\n", n); printf("CPU:\n"); printf("\tResult: %d\n", hostSum); printf("\tTime: %.3f ms \n", duration); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaMemcpy(d_data, data2, bytes, cudaMemcpyHostToDevice); dim3 blockSize(512); dim3 gridSize((n - 1) / blockSize.x + 1); cudaEventRecord(start); reduceOnDevice<<<gridSize, blockSize>>>(d_data, d_sum, n); cudaEventRecord(stop); cudaMemcpy(sum, d_sum, bytes, cudaMemcpyDeviceToHost); cudaEventSynchronize(stop); float milliseconds = 0; cudaEventElapsedTime(&milliseconds, start, stop); int gpu_sum = 0; for (int k = 0; k < gridSize.x; k++) gpu_sum += sum[k]; printf("GPU:\n"); printf("\tResult: %d\n", gpu_sum); printf("\tTime: %.3f ms \n", milliseconds); printf("Host/Device: %.3f\n", duration/milliseconds); printf("___________________________________\n"); cudaFree(d_data); cudaFree(d_sum); free(data); free(sum); } return 0; } |
Các bạn chú ý, chúng ta cần phải đợi có kết quả trước để dùng tính tổng cho kết quả sau nên chúng ta sẽ phải đồng bộ bằng cách dùng __syncthreads()
Ở bài viết sau, chúng ta sẽ tối ưu lại việc sử dụng thread cũng như block một cách hiệu quả để tránh bị phân kỳ (một trong những vấn đề khi lập trình song song trên GPU với CUDA)
Buy me a coffee
Leave a Reply