Programming, Network, Office, Hardware, Software, Design, Photography.. I am just a GEEK
Bài viết tiếp theo trong series về Lập trình song song trên GPU bằng CUDA.
Ở bài viết trước mình đã cùng nhau thực hiện Cộng 2 vector, bài viết này cũng là 1 trong 3 bài cơ bản để chúng ta tìm hiểu cách song song hóa các thuật toán cũng như các chương trình của chúng ta sau này.
Chúng ta sẽ cài đặt thuật toán Nhân 2 ma trận 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 ma trận: 100 x 100
Phát sinh ngẫu nhiên giá trị của các ma trận đầ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
Thuật toán
|
1 2 3 4 5 6 7 8 9 10 11 12 |
for (int iy = 0; iy < ny; iy++) { for (int ix = 0; ix < nx; ix++) { float sum = 0; for (int k = 0; k < nx; k++) { sum += in1[iy * nx + k] * in2[k * nx + ix]; } out[iy * nx + ix] = sum; } } |
Cách thực hiện song song bằng CUDA cũng giống với mảng 1 chiều, ở đây chúng ta sử dụng mảng 2 chiều
|
1 2 3 4 5 6 7 8 9 10 |
int ix = blockIdx.x*blockDim.x+threadIdx.x; int iy = blockIdx.y*blockDim.y+threadIdx.y; if (ix < nx && iy < ny) { float sum = 0; for (int k = 0; k < nx; k++) { sum += d_in1[iy * nx + k] * d_in2[k * nx + ix]; } d_out[iy * nx + ix] = sum; } |
Các bạn có thể tham khảo lại bài viết trước của mình cách để chạy CUDA trên Google Colab nhé.
|
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 |
#include <stdio.h> #include <stdlib.h> #include <math.h> #include <time.h> void multiplyMatOnHost(float *in1, float *in2, float *out, int nx, int ny) { for (int iy = 0; iy < ny; iy++) { for (int ix = 0; ix < nx; ix++) { float sum = 0; for (int k = 0; k < nx; k++) { sum += in1[iy * nx + k] * in2[k * nx + ix]; } out[iy * nx + ix] = sum; } } } __global__ void multiplyMatOnDevice(float *d_in1, float *d_in2, float *d_out, int nx, int ny) { int ix = blockIdx.x*blockDim.x+threadIdx.x; int iy = blockIdx.y*blockDim.y+threadIdx.y; if (ix < nx && iy < ny) { float sum = 0; for (int k = 0; k < nx; k++) { sum += d_in1[iy * nx + k] * d_in2[k * nx + ix]; } d_out[iy * nx + ix] = sum; } } int main(int argc, char* argv[]) { int sizeArray[10][2] = {{32, 32}, {16, 32}, {32, 16}, {16, 16}, {32, 1}, {64, 1}, {128, 1}, {256, 1}, {512, 1}, {1024, 1}}; for(int i = 0; i < 10; i++) { int nx = 100, ny = 100; float *in1, *in2, *out; float *d_in1, *d_in2, *d_out; size_t bytes = nx * ny * sizeof(float); in1 = (float*)malloc(bytes); in2 = (float*)malloc(bytes); out = (float*)malloc(bytes); cudaMalloc(&d_in1, bytes); cudaMalloc(&d_in2, bytes); cudaMalloc(&d_out, bytes); for (int m = 0; m < ny; m++){ for (int n = 0; n < nx; n++){ in1[m*nx+n] = sin(n)*sin(n); in2[m*nx+n] = cos(n)*cos(n); } } clock_t begin, end; begin = clock(); multiplyMatOnHost(in1, in2, out, nx, ny); end = clock(); double duration = double(end-begin)/CLOCKS_PER_SEC * 1000; printf("(Block size %dx%d):\n", sizeArray[i][0], sizeArray[i][1]); printf("CPU:\n"); printf("\tTime: %.3f ms \n", duration); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaMemcpy(d_in1, in1, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_in2, in2, bytes, cudaMemcpyHostToDevice); dim3 blockSize(sizeArray[i][0], sizeArray[i][1]); dim3 gridSize((nx - 1) / blockSize.x + 1, (ny - 1) / blockSize.y + 1); cudaEventRecord(start); multiplyMatOnDevice<<<gridSize, blockSize>>>(d_in1, d_in2, d_out, nx, ny); cudaEventRecord(stop); cudaMemcpy(out, d_out, bytes, cudaMemcpyDeviceToHost); cudaEventSynchronize(stop); float milliseconds = 0; cudaEventElapsedTime(&milliseconds, start, stop); printf("GPU:\n"); printf("\tTime: %.3f ms \n", milliseconds); printf("Host/Device: %.3f\n", duration/milliseconds); printf("___________________________________\n"); cudaFree(d_in1); cudaFree(d_in2); cudaFree(d_out); free(in1); free(in2); free(out); } return 0; } |
Buy me a coffee
Leave a Reply