11.更多的并行化 |
前面提到了 block。在之前介绍呼叫 CUDA 函式时,也有提到 "block 数目" 这个参数。到目前为止,我们都只使用一个 block。究竟 block 是什么呢?
在 CUDA 中,thread 是可以分组的,也就是 block。一个 block 中的 thread,具有一个共享的 shared memory,也可以进行同步工作。不同 block 之间的 thread 则不行。在我们的程序中,其实不太需要进行 thread 的同步动作,因此我们可以使用多个 block 来进一步增加 thread 的数目。
首先,在 define DATA_SIZE 的地方,改成如下:
#define DATA_SIZE 1048576
#define BLOCK_NUM 32
#define THREAD_NUM 256
这表示我们会建立 32 个 blocks,每个 blocks 有 256 个 threads,总共有 32*256 = 8192 个 threads。
接着,我们把 kernel 部份改成:
__global__ static void sumOfSquares(int *num, int* result,
clock_t* time)
{
const int tid = threadIdx.x;
const int bid = blockIdx.x;
int sum = 0;
int i;
if(tid == 0) time[bid] = clock();
for(i = bid * THREAD_NUM + tid; i < DATA_SIZE;
i += BLOCK_NUM * THREAD_NUM) {
sum += num[i] * num[i];
}
result[bid * THREAD_NUM + tid] = sum;
if(tid == 0) time[bid + BLOCK_NUM] = clock();
}
blockIdx.x 和 threadIdx.x 一样是 CUDA 内建的变量,它表示的是目前的 block 编号。另外,注意到我们把计算时间的方式改成每个 block 都会记录开始时间及结束时间。
main 函式部份,修改成:
int* gpudata, *result;
clock_t* time;
cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE);
cudaMalloc((void**) &result,
sizeof(int) * THREAD_NUM * BLOCK_NUM);
cudaMalloc((void**) &time, sizeof(clock_t) * BLOCK_NUM * 2);
cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE,
cudaMemcpyHostToDevice);
sumOfSquares<<<BLOCK_NUM, THREAD_NUM, 0>>>(gpudata, result,
time);
int sum[THREAD_NUM * BLOCK_NUM];
clock_t time_used[BLOCK_NUM * 2];
cudaMemcpy(&sum, result, sizeof(int) * THREAD_NUM * BLOCK_NUM,
cudaMemcpyDeviceToHost);
cudaMemcpy(&time_used, time, sizeof(clock_t) * BLOCK_NUM * 2,
cudaMemcpyDeviceToHost);
cudaFree(gpudata);
cudaFree(result);
cudaFree(time);
int final_sum = 0;
for(int i = 0; i < THREAD_NUM * BLOCK_NUM; i++) {
final_sum += sum[i];
}
clock_t min_start, max_end;
min_start = time_used[0];
max_end = time_used[BLOCK_NUM];
for(int i = 1; i < BLOCK_NUM; i++) {
if(min_start > time_used[i])
min_start = time_used[i];
if(max_end < time_used[i + BLOCK_NUM])
max_end = time_used[i + BLOCK_NUM];
}
printf("sum: %d time: %d\n", final_sum, max_end - min_start);
基本上我们只是把 result 的大小变大,并修改计算时间的方式,把每个 block 最早的开始时间,和最晚的结束时间相减,取得总运行时间。
这个版本的程序,执行的时间减少很多,在 GeForce 8800GT 上只需要约 150K cycles,相当于 40GB/s 左右的带宽。不过,它在 CPU 上执行的部份,需要的时间加长了(因为 CPU 现在需要加总 8192 个数字)。为了避免这个问题,我们可以让每个 block 把自己的每个 thread 的计算结果进行加总。
12.Thread 的同步 |
前面提过,一个 block 内的 thread 可以有共享的内存,也可以进行同步。我们可以利用这一点,让每个 block 内的所有 thread 把自己计算的结果加总起来。把 kernel 改成如下:
__global__ static void sumOfSquares(int *num, int* result,
clock_t* time)
{
extern __shared__ int shared[];
const int tid = threadIdx.x;
const int bid = blockIdx.x;
int i;
if(tid == 0) time[bid] = clock();
shared[tid] = 0;
for(i = bid * THREAD_NUM + tid; i < DATA_SIZE;
i += BLOCK_NUM * THREAD_NUM) {
shared[tid] += num[i] * num[i];
}
__syncthreads();
if(tid == 0) {
for(i = 1; i < THREAD_NUM; i++) {
shared[0] += shared[i];
}
result[bid] = shared[0];
}
if(tid == 0) time[bid + BLOCK_NUM] = clock();
}
利用 __shared__ 声明的变量表示这是 shared memory,是一个 block 中每个 thread 都共享的内存。它会使用在 GPU 上的内存,所以存取的速度相当快,不需要担心 latency 的问题。
__syncthreads() 是一个 CUDA 的内部函数,表示 block 中所有的 thread 都要同步到这个点,才能继续执行。在我们的例子中,由于之后要把所有 thread 计算的结果进行加总,所以我们需要确定每个 thread 都已经把结果写到 shared[tid] 里面了。
接下来,把 main 函式的一部份改成:
int* gpudata, *result;
clock_t* time;
cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE);
cudaMalloc((void**) &result, sizeof(int) * BLOCK_NUM);
cudaMalloc((void**) &time, sizeof(clock_t) * BLOCK_NUM * 2);
cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE,
cudaMemcpyHostToDevice);
sumOfSquares<<<BLOCK_NUM, THREAD_NUM,
THREAD_NUM * sizeof(int)>>>(gpudata, result, time);
int sum[BLOCK_NUM];
clock_t time_used[BLOCK_NUM * 2];
cudaMemcpy(&sum, result, sizeof(int) * BLOCK_NUM,
cudaMemcpyDeviceToHost);
cudaMemcpy(&time_used, time, sizeof(clock_t) * BLOCK_NUM * 2,
cudaMemcpyDeviceToHost);
cudaFree(gpudata);
cudaFree(result);
cudaFree(time);
int final_sum = 0;
for(int i = 0; i < BLOCK_NUM; i++) {
final_sum += sum[i];
}
可以注意到,现在 CPU 只需要加总 BLOCK_NUM 也就是 32 个数字就可以了。
不过,这个程序由于在 GPU 上多做了一些动作,所以它的效率会比较差一些。在 GeForce 8800GT 上,它需要约 164K cycles。
当然,效率会变差的一个原因是,在这一版的程序中,最后加总的工作,只由每个 block 的 thread 0 来进行,但这并不是最有效率的方法。理论上,把 256 个数字加总的动作,是可以并行化的。最常见的方法,是透过树状的加法:
把 kernel 改成如下:
__global__ static void sumOfSquares(int *num, int* result,
clock_t* time)
{
extern __shared__ int shared[];
const int tid = threadIdx.x;
const int bid = blockIdx.x;
int i;
int offset = 1, mask = 1;
if(tid == 0) time[bid] = clock();
shared[tid] = 0;
for(i = bid * THREAD_NUM + tid; i < DATA_SIZE;
i += BLOCK_NUM * THREAD_NUM) {
shared[tid] += num[i] * num[i];
}
__syncthreads();
while(offset < THREAD_NUM) {
if((tid & mask) == 0) {
shared[tid] += shared[tid + offset];
}
offset += offset;
mask = offset + mask;
__syncthreads();
}
if(tid == 0) {
result[bid] = shared[0];
time[bid + BLOCK_NUM] = clock();
}
}
后面的 while 循环就是进行树状加法。main 函式则不需要修改。
这一版的程序,在 GeForce 8800GT 上执行需要的时间,大约是 140K cycles(相当于约 43GB/s),比完全不在 GPU 上进行加总的版本还快!这是因为,在完全不在 GPU 上进行加总的版本,写入到 global memory 的数据数量很大(8192 个数字),也对效率会有影响。所以,这一版程序不但在 CPU 上的运算需求降低,在 GPU 上也能跑的更快。