Mã dưới đây báo cáo một ví dụ về Chỉ thị mức song song (ILP).
Chức năng __global__ trong ví dụ này chỉ thực hiện một phép gán giữa hai mảng. Đối với trường hợp ILP=1, chúng tôi có nhiều chủ đề như số lượng các phần tử mảng N, sao cho mỗi luồng thực hiện một nhiệm vụ duy nhất. Ngược lại, đối với trường hợp ILP=2, chúng tôi có một số chủ đề N/2 mỗi luồng xử lý 2 phần tử. Nói chung, đối với trường hợp ILP=k, chúng tôi có một số lượng N/k mỗi chủ đề xử lý k yếu tố.
Bên cạnh mã, bên dưới tôi cũng đang báo cáo thời gian, được thực hiện trên NVIDIA GT920M (kiến trúc Kepler), với các giá trị khác nhau của N và ILP. Vì nó có thể được nhìn thấy:
- cho các giá trị lớn
N, băng thông bộ nhớ gần với một tối đa cho thẻ GT920M, cụ thể là, 14.4GB/s, đạt được;
- cho bất kỳ cố định
N, thay đổi giá trị ILP không thay đổi hiệu suất.
Liên quan đến điểm 2., tôi cũng đã thử nghiệm cùng một mã trên Maxwell và quan sát cùng một hành vi (không thay đổi hiệu suất so với ILP).Để thay đổi hiệu suất so với ILP, vui lòng xem câu trả lời cho số The efficiency and performance of ILP for the NVIDIA Kepler architecture báo cáo cũng kiểm tra kiến trúc Fermi.
Tốc độ bộ nhớ đã được tính toán theo công thức sau:
(2.f * 4.f * N * numITER)/(1e9 * timeTotal * 1e-3)
nơi
4.f * N * numITER
là số đọc hoặc viết,
2.f * 4.f * N * numITER
là số đọc VÀ viết,
timeTotal * 1e-3
là thời gian ở seconds (timeTotal là ở ms).
BỘ LUẬT
// --- GT920m - 14.4 GB/s
// http://gpuboss.com/gpus/GeForce-GTX-280M-vs-GeForce-920M
#include<stdio.h>
#include<iostream>
#include "Utilities.cuh"
#include "TimingGPU.cuh"
#define BLOCKSIZE 32
#define DEBUG
/****************************************/
/* INSTRUCTION LEVEL PARALLELISM KERNEL */
/****************************************/
__global__ void ILPKernel(const int * __restrict__ d_a, int * __restrict__ d_b, const int ILP, const int N) {
const int tid = threadIdx.x + blockIdx.x * blockDim.x * ILP;
if (tid >= N) return;
for (int j = 0; j < ILP; j++) d_b[tid + j * blockDim.x] = d_a[tid + j * blockDim.x];
}
/********/
/* MAIN */
/********/
int main() {
//const int N = 8192;
const int N = 524288 * 32;
//const int N = 1048576;
//const int N = 262144;
//const int N = 2048;
const int numITER = 100;
const int ILP = 16;
TimingGPU timerGPU;
int *h_a = (int *)malloc(N * sizeof(int));
int *h_b = (int *)malloc(N * sizeof(int));
for (int i = 0; i<N; i++) {
h_a[i] = 2;
h_b[i] = 1;
}
int *d_a; gpuErrchk(cudaMalloc(&d_a, N * sizeof(int)));
int *d_b; gpuErrchk(cudaMalloc(&d_b, N * sizeof(int)));
gpuErrchk(cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice));
gpuErrchk(cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice));
/**************/
/* ILP KERNEL */
/**************/
float timeTotal = 0.f;
for (int k = 0; k < numITER; k++) {
timerGPU.StartCounter();
ILPKernel << <iDivUp(N/ILP, BLOCKSIZE), BLOCKSIZE >> >(d_a, d_b, ILP, N);
#ifdef DEBUG
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
#endif
timeTotal = timeTotal + timerGPU.GetCounter();
}
printf("Bandwidth = %f GB/s; Num blocks = %d\n", (2.f * 4.f * N * numITER)/(1e6 * timeTotal), iDivUp(N/ILP, BLOCKSIZE));
gpuErrchk(cudaMemcpy(h_b, d_b, N * sizeof(int), cudaMemcpyDeviceToHost));
for (int i = 0; i < N; i++) if (h_a[i] != h_b[i]) { printf("Error at i = %i for kernel0! Host = %i; Device = %i\n", i, h_a[i], h_b[i]); return 1; }
return 0;
}
THI
GT 920M
N = 512 - ILP = 1 - BLOCKSIZE = 512 (1 block - each block processes 512 elements) - Bandwidth = 0.092 GB/s
N = 1024 - ILP = 1 - BLOCKSIZE = 512 (2 blocks - each block processes 512 elements) - Bandwidth = 0.15 GB/s
N = 2048 - ILP = 1 - BLOCKSIZE = 512 (4 blocks - each block processes 512 elements) - Bandwidth = 0.37 GB/s
N = 2048 - ILP = 2 - BLOCKSIZE = 256 (4 blocks - each block processes 512 elements) - Bandwidth = 0.36 GB/s
N = 2048 - ILP = 4 - BLOCKSIZE = 128 (4 blocks - each block processes 512 elements) - Bandwidth = 0.35 GB/s
N = 2048 - ILP = 8 - BLOCKSIZE = 64 (4 blocks - each block processes 512 elements) - Bandwidth = 0.26 GB/s
N = 2048 - ILP = 16 - BLOCKSIZE = 32 (4 blocks - each block processes 512 elements) - Bandwidth = 0.31 GB/s
N = 4096 - ILP = 1 - BLOCKSIZE = 512 (8 blocks - each block processes 512 elements) - Bandwidth = 0.53 GB/s
N = 4096 - ILP = 2 - BLOCKSIZE = 256 (8 blocks - each block processes 512 elements) - Bandwidth = 0.61 GB/s
N = 4096 - ILP = 4 - BLOCKSIZE = 128 (8 blocks - each block processes 512 elements) - Bandwidth = 0.74 GB/s
N = 4096 - ILP = 8 - BLOCKSIZE = 64 (8 blocks - each block processes 512 elements) - Bandwidth = 0.74 GB/s
N = 4096 - ILP = 16 - BLOCKSIZE = 32 (8 blocks - each block processes 512 elements) - Bandwidth = 0.56 GB/s
N = 8192 - ILP = 1 - BLOCKSIZE = 512 (16 blocks - each block processes 512 elements) - Bandwidth = 1.4 GB/s
N = 8192 - ILP = 2 - BLOCKSIZE = 256 (16 blocks - each block processes 512 elements) - Bandwidth = 1.1 GB/s
N = 8192 - ILP = 4 - BLOCKSIZE = 128 (16 blocks - each block processes 512 elements) - Bandwidth = 1.5 GB/s
N = 8192 - ILP = 8 - BLOCKSIZE = 64 (16 blocks - each block processes 512 elements) - Bandwidth = 1.4 GB/s
N = 8192 - ILP = 16 - BLOCKSIZE = 32 (16 blocks - each block processes 512 elements) - Bandwidth = 1.3 GB/s
...
N = 16777216 - ILP = 1 - BLOCKSIZE = 512 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.9 GB/s
N = 16777216 - ILP = 2 - BLOCKSIZE = 256 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.8 GB/s
N = 16777216 - ILP = 4 - BLOCKSIZE = 128 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.8 GB/s
N = 16777216 - ILP = 8 - BLOCKSIZE = 64 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.7 GB/s
N = 16777216 - ILP = 16 - BLOCKSIZE = 32 (32768 blocks - each block processes 512 elements) - Bandwidth = 12.6 GB/s
Một out-of-trật tự xử lý không cần thiết để khai thác hướng dẫn mức xử lý song song. Một bộ xử lý theo thứ tự với thực thi superscalar cũng có thể có lợi. – njuffa