问题描述:将两组数据对应的元素两两相加,并将结果保存在第三个数组中。即矢量求和运算
目录
问题:索引取值范为0 到 N-1 ,下面两种函数add()有什么区别吗?
问题:既然GPU将运行kernel的N个副本,那如何在代码中知道当前正在运行是哪一个Block?
问题:为什么不是blockIdx?而是blockIdx.x?
1.基于CPU的矢量求和
代码分析:
void add(int* a,int* b,int* c){
int tid = 0; //CPU zero
while(tid < N){
c[tid] = a[tid] + b[tid];
tid += 1; //Only one CPU, increment by one
}
}
void add(int* a,int* b,int* c){
int i = 0;
for(i = 0; i < N ; i++ ){
c[i] = a[i] + b[i];
}
}
问题:索引取值范为0 到 N-1 ,下面两种函数add()有什么区别吗?
while循环虽然比较复杂,但是能够使代码在拥有多个CPU或者GPU核的系统上运行。
如下图:假如双核处理器,每次递增大小改为2,核1将偶数索引的元素相加,而核2将奇数索引的元素相加,即相当与每个CPU核上执行以下代码。
后面再实现这个代码,线程调度机制的实际运行情况灰常复杂
void add(int* a,int* b,int* c){
int tid = 0;
while(tid < N){
c[tid] = a[tid] + b[tid];
tid += 2;
}
}
void add(int* a,int* b,int* c){
int tid = 1;
while(tid < N){
c[tid] = a[tid] + b[tid];
tid += 2;
}
}
完整的代码实现及运行结果:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#define N 10
void add(int* a,int* b,int* c){
int tid = 0; //这是第0个CPU,因此索引从0开始
while(tid < N){
c[tid] = a[tid] + b[tid];
tid += 1; //由于只有一个CPU,因此每次递增1
}
}
int main(void){
int a[N],b[N],c[N];
//在CPU上为数组‘a’和‘b’赋值
for(int i =0;i<N;i++){
a[i] = -i;
b[i] = i * i;
}
add(a,b,c);
//显示结果
for(int i=0;i<N;i++){
printf("%d + %d = %d\n",a[i],b[i],c[i]);
}
return 0;
}
2.基于GPU的矢量求和
代码分析:
问题:<<<>>>里面参数式什么含义?
add<<<N,1>>>(dev_a,dev_b,dev_c);
N:device在执行kernel时,使用的Block数量。(后续会解释Block是什么)
即运行时创建kernel的N个副本,并以并行的方式来运行它们。
问题:既然GPU将运行kernel的N个副本,那如何在代码中知道当前正在运行是哪一个Block?
__global__ void add(int* a,int* b,int* c){
int tid = blockIdx.x; //计算索引处的数据
if(tid < N){
c[tid] = a[tid] + b[tid];
}
}
blockIdx是一个内置变量,在CUDA运行时中已经预先定义了这个变量,而且这个变量的名字就是变量的作用,变量中包含的值就是当前执行Device Code的Block的索引。
问题:为什么不是blockIdx?而是blockIdx.x?
因为CUDA支持二维的Block数组。对于二维空间计算问题,如矩阵数学运算或者图像处理,使用二维索引往往会带来很大的遍历,(后续使用再分析),避免将线性所以转化为矩形索引。
问题:进一步解释Block是如何定位kernel?
Block的集合称为Grid。在本程序中,我们使用的是一维的Grid,其中包含N个Block。每个Block的blockIdx.x值是不同的,如:第一个Block的blockIdx.x=0,第N个Block的block.x=N-1。
所以在本程序中,有N个Block,每个Block都运行相同的Device Code,但每个Block的blockIdx.x是不同的,所以当N个Block并行执行时,运行时用相应的Block索引来替换blockIdx.x。每个Block执行的Block Code如下。
__global__ void add(int* a,int* b,int* c){
int tid = 0;
if(tid < N){
c[tid] = a[tid] + b[tid];
}
}
//int tid = 1 ....int tid = N-1;
......
__global__ void add(int* a,int* b,int* c){
int tid = N-1;
if(tid < N){
c[tid] = a[tid] + b[tid];
}
}
问题:为什么要判断tid是否小于N?
在我们调整参数<<<>>>,tid总是小于N的,因为我们在kernel中都是这么假设的。但是我们任然有理由怀疑可能会出现非法的内存访问,这将会造成一种糟糕的情况。
问题:Block的数量是否可以任意设置?
在启动Block数组时,数组每一维的最大数量都不能超过65535,这是一种硬件限制,如果启动的Block数量超过了这个限制,那么程序将运行失败。
完整的代码实现及运行结果:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#define N 10
__global__ void add(int* a,int* b,int* c){
int tid = blockIdx.x; //计算索引处的数据
if(tid < N){
c[tid] = a[tid] + b[tid];
}
}
int main(void){
int a[N],b[N],c[N];
int *dev_a,*dev_b,*dev_c;
//在cpu上为数组‘a’和‘b’赋值
//为什么要在CPU上对输入数组赋值,其实没啥特殊原因?
//假设用于从硬盘上读取数据,或者是其他应用程序中的一个步骤,并且输入数组由其他算法生成
for (int i = 0; i<N; i++) {
a[i] = -i;
b[i] = i * i;
}
//1.在GPU上分配内存
cudaMalloc((void**)&dev_a, N * sizeof(int));
cudaMalloc((void**)&dev_b, N * sizeof(int));
cudaMalloc((void**)&dev_c, N * sizeof(int));
//2.将数组‘a’和‘b’复制到GPU
cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
//3.执行kernel
add<<<N,1>>>(dev_a,dev_b,dev_c);
//4.将数组‘c’从GPU复制到CPU
cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
//在CPU上显示结果
for (int i = 0; i<N; i++) {
printf("%d + %d = %d\n", a[i], b[i], c[i]);
}
//5.释放在GPU上分配的内存。目的避免内存泄露
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
return 0;
}
3.小结
本次学习笔记的总结,我们学到了如何告诉CUDA运行时在Block上并行执行程序,我们把在GPU上启动的Block集合称为一个Grid。Grid既可以是一维的Block集合,也可以是二维的Block集合。Kernel的每个副本都可以通过内置变量blockid来判断哪个Block正在执行它。同样也可以通过内置变量gridDim来获得Grid的大小。后续会用到