截止现在:将Block分解为Thread的目的只是为了解决Block数量的硬件限制。
问题:CUDA C的关键字__share__添加到变量声明中,这将使这个变量驻留在shared memory中,这样做的目的是什么?
CUDA C编辑器对shared memory中的变量与普通变量将分别采取不同的处理方式。对于GPU上启动的每个Block,CUDA C编译器都将创建该变量的一个副本。Block中的每个Thread都共享这块内存,但Thread却无法看到也不能修改其他Block的变量副本。这使得一个Block中的多个Thread能够在GPU上进行通信和协作。而且,shared memory驻留在物理GPU上,不是驻留在GPU之外的系统的memory中。因此,在访问shared memory时的延迟要远远低于访问普通缓冲区的延迟,使得shared memory像每个block的高速缓存或者中间结果暂存器那样高效。
问题:如果没有同步,那么将发生竞争条件(Race Condition),这种情况下,代码执行结果正确性将取决于硬件的不确定性。
目录
问题:为什么只有cacheIndex==0的线程执行这个保存操作?
问题:32个Block,并且每个Block包含256个线程,那么是否会造成线程过多的情况?
点积运算
描述:
每个Thread将两个对应的元素相乘,然后移动到下两个元素,由于最终的结果是所有乘积的总和,因此每个Thread还要保存它所计算的成绩的加和。
优化:归约(Reduction)算法
基本思想:对一输入数组执行某种计算,然后产生一个更小的结果数组。。每个Thread将cache[]中的两个值相加起来,然后将结果保存回cache[]。由于每个Thread都将两个值合并为一个值,那么在完成这个步骤后,得到的结果数量就是计算开始时数值数量的一半。在下一个步骤中,我们对这一半数值执行相同的操作。在将这种操作执行log2(threadsPerBlock)个步骤后,就能能得到cache[]中所有值的综合。
具体实现:第一步,去threadsPerBlock的一半作为i值,只有索引小于这个值的线程才会执行。只有当线程的索引小于i时,才可以把cache[]的两个数据项相加起来,因此将加法运算放在if条件里面,执行加法运算的线程将cache[]中线程索引位置上的值和线程索引加上i得到的位置上的值相加起来,并将结果保存回cache[]中线程索引位置上。
// 对于归约(reductions)运算来说, threadsPerBlock必须是2的整数倍
int i = blockDim.x / 2;
while (i != 0) {
if (cacheIndex < i)
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
i /= 2;
}
if (cacheIndex == 0)
[blockIdx.x] = cache[0];
问题:为什么只有cacheIndex==0的线程执行这个保存操作?
结束while()循环,每个Block都得到一个值,这个值位于cache[]的第一个元素中,并且就等于该block中两两元素乘积的加和。然后,我们将这个值保存到shared memory并结束Kernel。因为只有一个值写入到shared memory,因此只需要一个Thread来执行这个操作,当然,每个Thread都可以执行这个写入操作,但这么做将使得在写入单个值时带来不必要的内存通信量。
最终,得到数组c[],其中该数组的每个元素中都包含了某个Block计算得到的加和值,点击运算的最后一个步骤就是计算c[]中所有元素的和。
问题:为什么在尚未计算完成之前就返回到主机?
像GPU这种大规模的并行机器在执行最后的归约步骤时,通常会浪费计算资源,因为此时的数据集往往非常小,因此我们将执行控制返回给Host,并且由CPU来完成最后一个加法步骤,即将计算数组c[]中所有元素的和
if (cacheIndex == 0)
[blockIdx.x] = cache[0];
问题:32个Block,并且每个Block包含256个线程,那么是否会造成线程过多的情况?
如果有N个数据元素,那么通常只需要N个Thread来计算点积。这里的情况是,Thread数量应为ThreadsPerBlock的最小整数倍,并且要大于或者N。下面代码与矢量加法示例遇到情况类似。
const int N = 33*1024;
const int threadsPerBlock = 256;
const int blocksPerGrid =imin(32, (N + threadsPerBlock - 1) / threadsPerBlock);
问题:对__syncthreads的进一步研究
如果将代码修改为只等待那些需要写入共享内存的线程,是不是就能获得性能提升? 第二个_syncthreads,是因为在循环迭代中更新了shared memroy变量cache[],并且在循环的下一次迭代开始之前,需要确保当前迭代中所有线程的更新操作都已经完成了。
int myVar = 0;
if (threadIdx.x % 2)
myVar = threadIdx.x
当Thrad到达粗体的代码行时,当奇数索引的Thread执行这条指令时,偶数索引的Thread就不会执行任何操作。当某些Thread需要执行一条指令,而其他线程不需要执行时,这种情况称之为发散(Thread Divergence)。在正常的环境中,发散的分支只会使得某些线程处于空闲状态,而其他Thread将执行分支中的代码。
但在__syncthreads()中,Thread发散的造成的结果更加糟糕,CUDA架构将确保,除非Block中的每个Thread都执行了__syncthreads(),因此硬件将使这些线程等待等待,一直等,一直等,永久地等待下去。
因此,如果将__syncthreads()调用移入if()Block中,GPU将停止使用,那么任何cacheIndex大于或等于i的线程将永远都不能执行__syncthreads。这将使处理器挂起,因为GPU在等待某个永远都不会发生的事件。
int i = blockDim.x / 2;
while (i != 0) {
if (cacheIndex < i)
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
i /= 2;
}
int i = blockDim.x / 2;
while (i != 0) {
if (cacheIndex < i)
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
i /= 2;
}
完整程序及运行结果:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
#define imin(a,b) (a<b?a:b)
const int N = 33 * 1024;
const int threadsPerBlock = 256;
const int blocksPerGrid =
imin(32, (N + threadsPerBlock - 1) / threadsPerBlock);
//假设有32个block,每个block有256个thread
//tid = tid + 32 * 256
__global__ void dot(float *a, float *b, float *c) {
__shared__ float cache[threadsPerBlock];//Block中的每个Thread都能将它计算的临时结果保存到某个位置上
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int cacheIndex = threadIdx.x; //shared memory的偏移就等于thredIdx
float temp = 0;
while (tid < N) {
temp += a[tid] * b[tid];
tid += blockDim.x * gridDim.x;
}
//设置cache中相应位置上的值
cache[cacheIndex] = temp;
// 对Block中的Thread进行同步(synchronize)
__syncthreads();
// 对于归约(reductions)运算来说, threadsPerBlock必须是2的整数倍
int i = blockDim.x / 2;
while (i != 0) {
if (cacheIndex < i)
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
i /= 2;
}
if (cacheIndex == 0)
c[blockIdx.x] = cache[0];
}
int main(void) {
float *a, *b, c, *partial_c;
float *dev_a, *dev_b, *dev_partial_c;
//在CPU上分配内存
a = (float*)malloc(N * sizeof(float));
b = (float*)malloc(N * sizeof(float));
partial_c = (float*)malloc(blocksPerGrid * sizeof(float));
// 1.在GPU上分配内存
cudaMalloc((void**)&dev_a,
N * sizeof(float));
cudaMalloc((void**)&dev_b,
N * sizeof(float));
cudaMalloc((void**)&dev_partial_c,
blocksPerGrid * sizeof(float));
//填充CPU主机内存
for (int i = 0; i<N; i++) {
a[i] = i;
b[i] = i * 2;
}
// 2.将数组‘a’和‘b’复制到GPU上
cudaMemcpy(dev_a, a, N * sizeof(float),
cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, N * sizeof(float),
cudaMemcpyHostToDevice);
//3.执行kernel
dot << <blocksPerGrid, threadsPerBlock >> >(dev_a, dev_b,
dev_partial_c);
// 4.将数组‘c’从GPU复制到CPU
cudaMemcpy(partial_c, dev_partial_c,
blocksPerGrid * sizeof(float),
cudaMemcpyDeviceToHost);
// 在CPU上完成最终的求和工作
c = 0;
for (int i = 0; i<blocksPerGrid; i++) {
c += partial_c[i];
}
#define sum_squares(x) (x*(x+1)*(2*x+1)/6)
printf("Does GPU value %.6g = %.6g?\n", c,
2 * sum_squares((float)(N - 1)));
// free memory on the gpu side
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_partial_c);
// free memory on the cpu side
free(a);
free(b);
free(partial_c);
}
验证方式:
小结:
我们已经知道了如何将Block进一步分解为更小的并行执行单元,这种并行执行单元也成为Thread。由矢量相加示例,升级到了如何实现任意长度矢量的加法。我们还学习了归约运算的示例,以及如何通过shared memory 和 同步 来实现这个运算。这个示例说明了如何对GPU与CPU进行协作从而完成运算。最后,我们还给出了当忽略同步给应用程序造成的问题。