CUDA编程结构、存储管理、线程管理杂谈
CUDA编程结构
一个异构环境，通常有多个CPU多个GPU，他们都通过PCIe总线相互通信，也是通过PCIe总线分隔开的。所以要区分一下两种设备的内存：
1）主机：CPU及其内存
2）设备：GPU及其内存
这两个内存从硬件到软件都是隔离的（CUDA6.0 以后支持统一寻址），目前先不研究统一寻址，现在还是用内存来回拷贝的方法来编写调试程序，以巩固大家对两个内存隔离这个事实的理解。
一个完整的CUDA应用可能的执行顺序，如图10-9所示。
图10-9 一个完整的CUDA应用可能的执行顺序
从host的串行到调用核函数（核函数被调用后控制马上归还主机线程，也就是在第一个并行代码执行时，很有可能第二段host代码已经开始同步执行了）。
接下来的研究层次是：
1）内存
2）线程
3）核函数
①启动核函数
②编写核函数
③验证核函数
4）错误处理
内存管理
内存管理在传统串行程序是非常常见的，寄存器空间，栈空间内的内存由机器自己管理，堆空间由用户控制分配和释放，CUDA程序同样，只是CUDA提供的API可以分配管理设备上的内存，当然也可以用CDUA管理主机上的内存，主机上的传统标准库也能完成主机内存管理。
一些主机API和CUDA C的API的对比，见表10-1。
表10-1 一些主机API和CUDA C的API的对比
|
标准C函数 |
CUDA C 函数 |
说明 |
|
malloc |
cudaMalloc |
内存分配 |
|
memcpy |
cudaMemcpy |
内存复制 |
|
memset |
cudaMemset |
内存设置 |
|
free |
cudaFree |
释放内存 |
图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值就不变了。他们是有区别的！这一点必须要注意。 下面有一段代码，块的索引和维度： /* *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; } 标签：存储管理,res,grid,CUDA,线程,block,内存 From： https://www.cnblogs.com/wujianming-110117/p/18386441