Programming, Network, Office, Hardware, Software, Design, Photography.. I am just a GEEK
Phần tiếp theo của bài viết Cộng các phần tử trong mảng – Lập trình song song trên GPU bằng CUDA
Bài viết này, chúng ta sẽ cải tiến lại bài toán tính tổng các phần tử trong mảng thực thi song song trên GPU bằng CUDA.
Kiến trúc phần cứng cơ bản của GPU
– GPU bao gồm các SM (Streaming Multiprocessor) – bộ xử lý đa luồng
- Mỗi SM bao gồm các SP (Streaming Processor) – bộ xử lý luồng (còn gọi là CUDA core)
– “Compute capability” = SM version
CUDA “ảo hóa” (virtualize) kiến trúc phần cứng của GPU
- Block = bộ-xử-lý-đa-luồng “ảo”
- Thread = bộ-xử-lý-luồng “ảo”
Khi host gọi hàm kernel, hệ thống sẽ tạo ra một grid gồm các block và mỗi block (bộ xử lý đa luồng “ảo”) sẽ được phân vào một SM (bộ xử lý đa luồng thật) để thực thi
- Mỗi SM có thể chứa nhiều hơn một block để thực thi
- Các block chưa được thực thi sẽ được đưa vào một hàng đợi
- Khi có một block được thực thi xong, hệ thống sẽ lấy một block chưa được thực thi ở hàng đợi và đưa vào thực thi
Hiện tượng phân kỳ warp
Trong SM, với mỗi block, hệ thống không quản lý và thực thi riêng lẻ từng thread mà làm theo các nhóm 32 thread – gọi là warp
Cách làm này được gọi là SIMT (Single Instruction Multiple Thread) – một câu lệnh được thực thi đồng thời cho tất các thread trong warp (mỗi thread có dữ liệu riêng của mình)
Warp bị phân kỳ (warp divergence)
- Nếu các thread trong warp không thể thực thi cùng một câu lệnh
Ví dụ: câu lệnh rẽ nhánh, vòng lặp
Tính tổng các phần tử trong mảng
Chúng ta cùng xem lại hàm reduceOnDevice của bài viết Cộng các phần tử trong mảng – Lập trình song song trên GPU bằng CUDA. Ta thấy wrap sẽ bị phân kỳ do có những thread lẻ sẽ không thực thi, vậy ta sẽ sửa lại 1 chút.
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 |
__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) { int index = 2 * stride * tid; if(index < blockDim.x) { idata[index] += idata[index + stride]; } __syncthreads(); } if(tid == 0) d_sum[blockIdx.x] = idata[0]; } |
Vậy chúng ta có thể tối ưu hóa sử dụng phần cứng của GPU hơn không?
Phiên bản Interleaved Pairs
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 |
__global__ void reduce1(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 = blockDim.x / 2; stride > 0; stride >>= 1) { if(tid < stride) { idata[tid] += idata[tid + stride]; } __syncthreads(); } if(tid == 0) d_sum[blockIdx.x] = idata[0]; } |
Phiên bản Unrolling
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 |
__global__ void reduce2(int *d_data, int *d_sum, int n) { int tid = threadIdx.x; int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x; int *idata = d_data + blockIdx.x * blockDim.x * 2; if(idx + blockDim.x < n) { d_data[idx] += d_data[idx + blockDim.x]; } __syncthreads(); for(int stride = blockDim.x / 2; stride > 0; stride >>= 1) { if(tid < stride) { idata[tid] += idata[tid + stride]; } __syncthreads(); } if(tid == 0) d_sum[blockIdx.x] = idata[0]; } |
Phiên bản Unrolled Warps
|
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 |
__global__ void reduce3 (int *d_data, int *d_sum, int n) { int tid = threadIdx.x; int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x; int *idata = d_data + blockIdx.x * blockDim.x * 8; if (idx + 7 * blockDim.x < n) { int a1 = d_data[idx]; int a2 = d_data[idx + blockDim.x]; int a3 = d_data[idx + 2 * blockDim.x]; int a4 = d_data[idx + 3 * blockDim.x]; int b1 = d_data[idx + 4 * blockDim.x]; int b2 = d_data[idx + 5 * blockDim.x]; int b3 = d_data[idx + 6 * blockDim.x]; int b4 = d_data[idx + 7 * blockDim.x]; d_data[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4; } __syncthreads(); for (int stride = blockDim.x / 2; stride > 32; stride >>= 1) { if (tid < stride) { idata[tid] += idata[tid + stride]; } __syncthreads(); } if (tid < 32) { volatile int *vmem = idata; vmem[tid] += vmem[tid + 32]; vmem[tid] += vmem[tid + 16]; vmem[tid] += vmem[tid + 8]; vmem[tid] += vmem[tid + 4]; vmem[tid] += vmem[tid + 2]; vmem[tid] += vmem[tid + 1]; } if (tid == 0) d_sum[blockIdx.x] = idata[0]; } |
Phiên bản Complete Unrolling
|
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 |
__global__ void reduce4 (int *d_data, int *d_sum, int n) { int tid = threadIdx.x; int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x; int *idata = d_data + blockIdx.x * blockDim.x * 8; if (idx + 7 * blockDim.x < n) { int a1 = d_data[idx]; int a2 = d_data[idx + blockDim.x]; int a3 = d_data[idx + 2 * blockDim.x]; int a4 = d_data[idx + 3 * blockDim.x]; int b1 = d_data[idx + 4 * blockDim.x]; int b2 = d_data[idx + 5 * blockDim.x]; int b3 = d_data[idx + 6 * blockDim.x]; int b4 = d_data[idx + 7 * blockDim.x]; d_data[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4; } __syncthreads(); if (blockDim.x>=1024 && tid < 512) idata[tid] += idata[tid + 512]; __syncthreads(); if (blockDim.x>=512 && tid < 256) idata[tid] += idata[tid + 256]; __syncthreads(); if (blockDim.x>=256 && tid < 128) idata[tid] += idata[tid + 128]; __syncthreads(); if (blockDim.x>=128 && tid < 64) idata[tid] += idata[tid + 64]; __syncthreads(); if (tid < 32) { volatile int *vsmem = idata; vsmem[tid] += vsmem[tid + 32]; vsmem[tid] += vsmem[tid + 16]; vsmem[tid] += vsmem[tid + 8]; vsmem[tid] += vsmem[tid + 4]; vsmem[tid] += vsmem[tid + 2]; vsmem[tid] += vsmem[tid + 1]; } if (tid == 0) d_sum[blockIdx.x] = idata[0]; } |
Buy me a coffee
Leave a Reply