一、CUDA编程模型概述

CUDA编程模型为应用和硬件设备之间的桥梁，所以CUDA C是编译型语言，不是解释型语言。

OpenCL就有点类似于解释型语言，通过编译器和链接，给操作系统执行（操作系统包括GPU在内的系统）

编译型语言：代码->编译->运行，编译成二进制，编译时报错

解释型语言：一边解释代码一边运行，运行到错误才报错

特点：

编译型语言：和硬件贴近，编译好速度快

解释型语言：不依赖硬件，更容易移植

上图，其中Communication Abstraction是编程模型和编译器，库函数之间的分界线。

编程模型：

模型，即共性。编程最重要的共性就是，程序设计时，代码的抽象方式、组织方式或复用方式。

也可以理解为，我们要用到的语法，内存结构，线程结构等这些我们写程序时我们自己控制的部分，这些部分控制了异构计算设备的工作模式，都是属于编程模型。

GPU中大致可以分为：

• 核函数
• 内存管理
• 线程管理
• 流

以上这些理论同时也适用于其他非CPU+GPU异构的组合。

从宏观上我们可以从以下几个环节完成CUDA应用开发：

• 领域层
• 逻辑层
• 硬件层

二、CUDA编程结构

一个异构环境，通常有多个CPU多个GPU，他们都通过PCIe总线相互通信，也是通过PCIe总线分隔开的。所以我们要区分一下两种设备的内存：

• 主机：CPU及其内存
• 设备：GPU及其内存

注意这两个内存从硬件到软件都是隔离的（CUDA6.0 以后支持统一寻址），我们目前先不研究统一寻址，我们现在还是用内存来回拷贝的方法来编写调试程序，以巩固大家对两个内存隔离这个事实的理解

从host的串行到调用核函数（核函数被调用后控制马上归还主机线程，也就是在第一个并行代码执行时，很有可能第二段host代码已经开始同步执行了）。

我们接下来的研究层次是：

• 内存
• 线程
• 核函数
  • 启动核函数
  • 编写核函数
  • 验证核函数
• 错误处理

三、内存管理

• 内存管理在传统串行程序是非常常见的，寄存器空间，栈空间内的内存由机器自己管理，堆空间由用户控制分配和释放，CUDA程序同样，只是CUDA提供的API可以分配管理设备上的内存，当然也可以用CDUA管理主机上的内存，主机上的传统标准库也能完成主机内存管理。
  下面表格有一些主机API和CUDA C的API的对比：

我们先研究最关键的一步，这一步要走总线的（郭德纲：我到底能不能走二环）

这个函数是内存拷贝过程，可以完成以下几种过程（cudaMemcpyKind kind）

• cudaMemcpyHostToHost
• cudaMemcpyHostToDevice
• cudaMemcpyDeviceToHost
• cudaMemcpyDeviceToDevice

这四个过程的方向可以清楚的从字面上看出来，这里就不废话了，如果函数执行成功，则会返回 cudaSuccess 否则返回 cudaErrorMemoryAllocation

使用下面这个指令可以吧上面的错误代码翻译成详细信息：

内存是分层次的，下图可以简单地描述，但是不够准确，后面我们会详细介绍每一个具体的环节：

共享内存（shared Memory）和全局内存（global Memory）后面我们会特别详细深入的研究，这里我们来个例子，两个向量的加法：

示例代码：

/* https://github.com/Tony-Tan/CUDA_Freshman 3_sum_arrays */ #include #include #include “freshman.h” void sumArrays(float a,float b,float * res,const int size)

{

for(int i=0;i4) { 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(floata,floatb,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;

// 分配设备端的内存空间，为了区分设备和主机端内存，我们可以给变量加后缀或者前缀h表示host，d表示device

// 一个经常会发生的错误就是混用设备和主机的内存地址！！

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<<>>(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;

}

四、线程管理

dio.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<<>>();

cudaDeviceReset();

return 0;

}

可以运行得到不同线程分解方式

此处有图，明天补上！

接下来这段代码是检查网格和块的大小的：

/*

*2_grid_block

*/

include

include

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;

}

这里也有图，明天补上

网格和块的维度存在几个限制因素，块大小主要与可利用的计算资源有关，如寄存器共享内存。

分成网格和块的方式可以使得我们的CUDA程序可以在任意的设备上执行。