CUDA编程结构、存储管理、线程管理杂谈
CUDA编程结构
一个异构环境，通常有多个CPU多个GPU，他们都通过PCIe总线相互通信，也是通过PCIe总线分隔开的。所以要区分一下两种设备的内存：
1）主机：CPU及其内存
2）设备：GPU及其内存
这两个内存从硬件到软件都是隔离的（CUDA6.0 以后支持统一寻址），目前先不研究统一寻址，现在还是用内存来回拷贝的方法来编写调试程序，以巩固大家对两个内存隔离这个事实的理解。
一个完整的CUDA应用可能的执行顺序，如图10-9所示。
图10-9 一个完整的CUDA应用可能的执行顺序
从host的串行到调用核函数（核函数被调用后控制马上归还主机线程，也就是在第一个并行代码执行时，很有可能第二段host代码已经开始同步执行了）。
从host的串行到调用核函数（核函数被调用后控制马上归还主机线程，也就是在第一个并行代码执行时，很有可能第二段host代码已经开始同步执行了）。
接下来的研究层次是：
1）内存
2）线程
3）核函数
①启动核函数
②编写核函数
③验证核函数
4）错误处理
内存管理
内存管理在传统串行程序是非常常见的，寄存器空间，栈空间内的内存由机器自己管理，堆空间由用户控制分配和释放，CUDA程序同样，只是CUDA提供的API可以分配管理设备上的内存，当然也可以用CDUA管理主机上的内存，主机上的传统标准库也能完成主机内存管理。
一些主机API和CUDA C的API的对比，见表10-1。
一些主机API和CUDA C的API的对比，见表10-1。
表10-1 一些主机API和CUDA C的API的对比
|
标准C函数 |
CUDA C 函数 |
说明 |
|
malloc |
cudaMalloc |
内存分配 |
|
memcpy |
cudaMemcpy |
内存复制 |
|
memset |
cudaMemset |
内存设置 |
|
free |
cudaFree |
释放内存 |
先研究最关键的一步，这一步要走总线的。
cudaError_t cudaMemcpy(void * dst,const void * src,size_t count,
cudaMemcpyKind kind)
这个函数是内存拷贝过程，可以完成以下几种过程（cudaMemcpyKind kind）
1）cudaMemcpyHostToHost
2）cudaMemcpyHostToDevice
3）cudaMemcpyDeviceToHost
4）cudaMemcpyDeviceToDevice
这四个过程的方向可以清楚的从字面上看出来，这里就不废话了，如果函数执行成功，则会返回 cudaSuccess 否则返回 cudaErrorMemoryAllocation
使用下面这个指令可以吧上面的错误代码翻译成详细信息：
char* cudaGetErrorString(cudaError_t error)
内存是分层次的，可以简单地描述，但是不够准确，后面会详细介绍每一个具体的环节，如图10-10所示。
图10-10 CUDA内存是分层次的
共享内存（shared Memory）和全局内存（global Memory）后面会特别详细深入的研究，这里来个例子，两个向量的加法：
/*
* 3_sum_arrays
*/
#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.h"
void sumArrays(float * a,float * b,float *
res,const int size)
{
for(int
i=0;i<size;i+=4)
{
res[i]=a[i]+b[i];
res[i+1]=a[i+1]+b[i+1];
res[i+2]=a[i+2]+b[i+2];
res[i+3]=a[i+3]+b[i+3];
}
}
__global__ void
sumArraysGPU(float*a,float*b,float*res)
{
int
i=threadIdx.x;
res[i]=a[i]+b[i];
}
int main(int argc,char **argv)
{
int
dev = 0;
cudaSetDevice(dev);
int
nElem=32;
printf("Vector size:%d\n",nElem);
int
nByte=sizeof(float)*nElem;
float *a_h=(float*)malloc(nByte);
float *b_h=(float*)malloc(nByte);
float *res_h=(float*)malloc(nByte);
float *res_from_gpu_h=(float*)malloc(nByte);
memset(res_h,0,nByte);
memset(res_from_gpu_h,0,nByte);
float *a_d,*b_d,*res_d;
CHECK(cudaMalloc((float**)&a_d,nByte));
CHECK(cudaMalloc((float**)&b_d,nByte));
CHECK(cudaMalloc((float**)&res_d,nByte));
initialData(a_h,nElem);
initialData(b_h,nElem);
CHECK(cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice));
dim3
block(nElem);
dim3
grid(nElem/block.x);
sumArraysGPU<<<grid,block>>>(a_d,b_d,res_d);
printf("Execution
configuration<<<%d,%d>>>\n",block.x,grid.x);
CHECK(cudaMemcpy(res_from_gpu_h,res_d,nByte,cudaMemcpyDeviceToHost));
sumArrays(a_h,b_h,res_h,nElem);
checkResult(res_h,res_from_gpu_h,nElem);
cudaFree(a_d);
cudaFree(b_d);
cudaFree(res_d);
free(a_h);
free(b_h);
free(res_h);
free(res_from_gpu_h);
return 0;
}
然后使用nvcc编译程序（代码库用cmake管理工程，更方便）。
解释下内存管理部分的代码：
cudaMalloc((float**)&a_d,nByte);
分配设备端的内存空间，为了区分设备和主机端内存，可以给变量加后缀或者前缀h_表示host，d_表示device
一个经常会发生的错误就是混用设备和主机的内存地址！！
线程管理
当内核函数开始执行，如何组织GPU的线程就变成了最主要的问题了，必须明确，一个核函数只能有一个grid，一个grid可以有很多个块，每个块可以有很多的线程，这种分层的组织结构使得并行过程更加自如灵活，如图10-11所示。
图10-11 CUDA线程管理示例
一个线程块block中的线程，可以完成下述协作：
1）同步
2）共享内存
不同块内线程不能相互影响！他们是物理隔离的！
接下来就是给每个线程一个编号了，知道每个线程都执行同样的一段串行代码，那么怎么让这段相同的代码对应不同的数据呢？首先第一步就是让这些线程彼此区分开，才能对应到相应从线程，使得这些线程也能区分自己的数据。如果线程本身没有任何标记，那么没办法确认其行为。
依靠下面两个内置结构体确定线程标号：
依靠下面两个内置结构体确定线程标号：
1）blockIdx（线程块在线程网格内的位置索引）
2）threadIdx（线程在线程块内的位置索引）
这里的Idx是index的缩写（之前一直以为是identity
x的缩写），这两个内置结构体基于 uint3 定义，包含三个无符号整数的结构，通过三个字段来指定：
1）blockIdx.x
2）blockIdx.y
3）blockIdx.z
4）threadIdx.x
5）threadIdx.y
6）threadIdx.z
上面这两个是坐标，当然要有同样对应的两个结构体来保存其范围，也就是blockIdx中三个字段的范围threadIdx中三个字段的范围：
1）blockDim
2）gridDim
他们是dim3类型(基于uint3定义的数据结构)的变量，也包含三个字段x,y,z.
1）blockDim.x
2）blockDim.y
3）blockDim.z
网格和块的维度一般是二维和三维的，也就是说一个网格通常被分成二维的块，而每个块常被分成三维的线程。
注意：dim3是手工定义的，主机端可见。uint3是设备端在执行的时候可见的，不可以在核函数运行时修改，初始化完成后uint3值就不变了。他们是有区别的！这一点必须要注意。
注意：dim3是手工定义的，主机端可见。uint3是设备端在执行的时候可见的，不可以在核函数运行时修改，初始化完成后uint3值就不变了。他们是有区别的！这一点必须要注意。
下面有一段代码，块的索引和维度：
/*
*1_check_dimension
*/
#include <cuda_runtime.h>
#include <stdio.h>
__global__ void checkIndex(void)
{
printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d)
blockDim:(%d,%d,%d)\
gridDim(%d,%d,%d)\n",threadIdx.x,threadIdx.y,threadIdx.z,
blockIdx.x,blockIdx.y,blockIdx.z,blockDim.x,blockDim.y,blockDim.z,
gridDim.x,gridDim.y,gridDim.z);
}
int main(int argc,char **argv)
{
int
nElem=6;
dim3
block(3);
dim3
grid((nElem+block.x-1)/block.x);
printf("grid.x %d grid.y %d grid.z %d\n",grid.x,grid.y,grid.z);
printf("block.x %d block.y %d block.z
%d\n",block.x,block.y,block.z);
checkIndex<<<grid,block>>>();
cudaDeviceReset();
return 0;
}
接下来这段代码是检查网格和块的大小的：
/*
*2_grid_block
*/
#include <cuda_runtime.h>
#include <stdio.h>
int main(int argc,char ** argv)
{
int
nElem=1024;
dim3
block(1024);
dim3
grid((nElem-1)/block.x+1);
printf("grid.x %d block.x %d\n",grid.x,block.x);
block.x=512;
grid.x=(nElem-1)/block.x+1;
printf("grid.x %d block.x %d\n",grid.x,block.x);
block.x=256;
grid.x=(nElem-1)/block.x+1;
printf("grid.x %d block.x %d\n",grid.x,block.x);
block.x=128;
grid.x=(nElem-1)/block.x+1;
printf("grid.x %d block.x %d\n",grid.x,block.x);
cudaDeviceReset();
return 0;
}
