同时，host与device之间进行通信，这样它们之间可以进行数据拷贝。典型的CUDA程序的执行的程序的流程为：

1. 分配host内存，并进行数据初始化；
2. 分配device内存（显存、共享内存），并从host端将数据拷贝到device端；
3. 调用CUDA的核函数在device函数上完成指定的运算；
4. 将device上的运算结果拷贝到host上；
5. 释放device和host上分配的内存。

由于CUDA编程模型实际上是异构编程模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词区别开host和device上的函数，主要的三个函数类型限定词如下：

• __global__: 在device端上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须为  void , 不支持可变参数，不能成为类成员函数。注意 __global__ 定义的kernel是异步的，这意味着host端不会等待kernel执行完就执行下一步；
• __device__: 在device端上执行，但仅可以从device中调用，不可以和  __global__ 同时用；
• __host__: 在host上执行，仅可以从host中调用，一般省略不写，不可以和  __global__同时用，但可以和  __device__，此时函数会在device和host都编译。

上面的流程中最重要的一个过程是调用CUDA的核函数来执行并行计算，kernel是CUDA中的一个重要的概念，kernel是在device上线程中并行执行的函数，在调用时需要用  <<<grid,block>>> 来指定kernel要执行的线程数量，在CUDA中，每个线程都要执行核函数，并且每个线程会分配一个唯一的 $threadID$ ,这个 $ID$ 值可以通过核函数的内置变量  thread Idx 来获得。

要深刻理解 $kernel$，必须要对 $kernel$ 的线程层次结构有一个清晰的认识。首先， $GPU$上很多并形化的轻量级线程。 $kernel$ 在device上执行时实际上是启动很多线程，一个 $kernel$ 所启动的所有线程称为 网格 $grid$ ，同一个网格的线程共享相同的全局内存空间，grid是线程结构的第一个层次，而网格又可以分为很多 线程块(block)，一个线程块里面包含很多线程，这是第二个层次。 为了编程方便， $grid$ 和 $block$ 都是定义为  dim3 类型的变量， dim3 可以看成是包含三个无符号整数 $(x,y,z)$ 成员的结构体变量，在定义时，缺失值初始化为1。因此，grid和block可以灵活地定义为1-dim，2-dim以及3-dim的结构，对于， $knernel$在定义调用时也必须通过执行配置  <<<grid,block>>>来指定kernel所使用的线程数及结构。

所以，为了方便编程，CUDA中使用了  dim3 类型（ dim3 是基于unit3定义的矢量类型，相当于由3个  unsigned int类型组成的结构体）的内建变量  threadIdx 和  blockIdx。这样，就可以使用一维、二维或三维的索引来标识线程，构成 一维、二维或三维线程块。使得线程组织形式对各种域（向量、矩阵，或者高维张量）中数据的划分变得直观、自然。

• 对于一维的block，线程的 $threadID$就是 $threadId.x$;
• 对于大小为 $(Dx,Dy)$的二维线程块block，线程的 $threadID$ 是  $(threadIdx.x+threadIdx.x\times Dx)$;
• 对于大小为 $(Dx,Dy,Dz)$的三维线程块block, 线程的 $threadID$是（ $threadIdx.x+threadIdx.y\times Dx+threadIdx.z\times Dx\times Dy$）;

另外，线程还有内置变量gridDim，用于获取网格块各个维度的大小。

此外，这里简单介绍一下CUDA的内存模块，如图所示。可以看到，每个线程有自己的私有本地内存（ $LocalMemory$）, 而每个线程块有包含共享内存（ $SharedMemory$）。还可以访问一些只读内存块：常用内存（ $ConstantMemory$）和纹理内存 （ $TextureMemory$）。内存结构涉及到程序优化，这里就过多讨论。

还有重要一点，你需要对 $GPU$的硬件实现有一个基本的认识。上面说到了 $kernel$的线程组织层次，那么一个 $kernel$ 实际上会启动很多线程，这些线程是逻辑上是并行的，但是在物理层也是无法却并不一定。这其实和CPU的多线程有类似之处，多线程如果没有多核支持，在物理层也无法实现并行的。但是好在 $GPU$ 存在很多CUDA核心，充分利用CUDA核心可以充分发挥GPU的并行计算能力。GPU硬件的一个核心组件是SM，前面已经说过，SM是Streaming Multiprocessor，SM的核心组件包括的CUDA核心、共享内存、寄存器等，SM可以并发的执行上，一个线程块只能在一个SM上被调度。SM一般可以调度多个线程块，这要看SM本身的能力。那么有可能一个kernel的各个线程块被分配多个SM，所以grid只是逻辑层，而SM才是执行的物理层。SM采用的是 SIMT(Single-Instruction, Multiple-Thread，单指令多线程)架构，基本的执行单元是线程束（warps)，线程束包含32个线程，这些线程同时执行相同的指令，但是每个线程都包含自己的指令地址计数器和寄存器状态，也有自己独立的执行路径。所以尽管线程束中的线程同时从同一程序地址执行，但是可能具有不同的行为，比如遇到了分支结构，一些线程可能进入这个分支，但是另外一些有可能不执行，它们只能死等，因为GPU规定线程束中所有线程在同一周期执行相同的指令，线程束分化会导致性能下降。当线程块被划分到某个SM上时，它将进一步划分为多个线程束，因为这才是SM的基本执行单元，但是一个SM同时并发的线程束数是有限的。这是因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程束中的线程分配独立的寄存器。所以SM的配置会影响其所支持的线程块和线程束并发数量。总之，就是网格和线程块只是逻辑划分，一个kernel的所有线程其实在物理层是不一定同时并发的。所以kernel的grid和block的配置不同，性能会出现差异，这点是要特别注意的。还有，由于SM的基本执行单元是包含32个线程的线程束，所以block大小一般要设置为32的倍数。

内存访问特性比较：

下面我将详细地介绍CUDA中各种内存管理函数的功能、参数和使用方法。

CUDA是一种用于异构并行计算的编程模型，经常需要在主机端（host）和设备端（Device）之间进行数据传输。这是因为CUDA核函数传入的必须是指向其中处理GPU显存的三个关键的API： cudaMalloc, cudaMemcpy和  cudaFree。

• cudaMalloc：

注意事项：

分配的内存

• cudaMemcpy ：

  其接口形式：	cudaError_t cudaMemcpy(void * dist, const void * src,size_t count,CudaMemcpyKind kind)
  函数功能：	将指向的内存区域的字节复制到指向的存储区域
  参数：	dist-目的存储地址；src -源内存地址；count-复制内存的字节数； kind-传输类型
  返回值：	cudaSuccess，cudaErrorInvalidValue，cudaErrorInvalidDevicePointer，cudaErrorInvalidMemcpyDirection

• cudaFree：

  其接口形式：	cudaError_t cudaFree(void * devPtr)
  函数功能：	释放由 指向的内存空间，该空间必须是之前调用cudaMalloc()或cudaMallocPitch()时返回过的。否则，或者如果cudaFree()之前已被调用过，则返回错误。如果 为 0，则不执行作。cudaFree() 在失败时返回cudaErrorInvalidDevicePointer。
  参数:	devPtr -设备指针指向内存释放
  返回值：	cudaSuccess，cudaErrorInvalidDevicePointer, cudaErrorInitialization

1.2 CUDA的第一个程序—矩阵赋值(Matrix Assign)

在本节通过一个矩阵赋值（matrix Assign）例子开始真正的CUDA程序实现，本例是在SDK中template程序的基础上修改得到的。 $template$ 是  $NVIDIA$ 公司提供的CUDA程序模板，也就是CUDA程序最基本的框架。要创建一个CUDA程序，可以把整个template文件复制一份。在一个CUDA程序中，基本的主机端代码主要完成以下的功能：

• 启动CUDA，使用多卡时应该时应该加上设备号，或使用 $cudaSetDevice\left(\right)$设备GPU设备；
• 为输入数据分配内存空间；
• 初始化输入数据；
• 为GPU分配内存，用于存放输入数据；
• 将内存中的输入数据拷贝到显存；
• 为GPU分配显存，用于存放输出数据；
• 调用device端的kernel进行计算，将结果写到显存中的对应区域；
• 为CPU分配内存，用于存放GPU传回来的输出数据；
• 将显存中的结果读取到内存；
• 释放内存和显存空间；
• 退出CUDA；

最简单的设备端代码主要完成以下功能：

• 从显存读取数据到GPU片内；

• 对数据进行处理；

• 将处理后的数据写回显存；

  其整个工程包含了三：

  （1）主程序文件CPU-Host端程序（example1main.cu）；

（2）GPU设备端函数的处理函数头文件（example_matrixassign_kernel.cuh）；

（3）GPU设备端函数的处理函数文件（example_matrixassign_kernel.cu）；

File1：主程序文件CPU-Host端程序（example1main.cu）；

#include<stdio.h> //系统头文件#include<stdlib.h>#include<string.h>#include<math.h>#include"cuda_runtime.h" //cuda项目头文件#include"device_launch_parameters.h"#include"example_matrixassign_kernel.cuh"  //核函数的数据的头文件void runTest(int argc, char** argv);int main(int argc,char** argv){
	runTest(argc,argv);
}void runTest(int gov.nme.zuzb.info argc, char** argv){	unsigned int num_blocks = 4;  //定义网格中的线程块数量
	unsigned int num_threads= 4;  //定义每个线程块中的线程数量
	unsigned int mem_size = sizeof(float) * num_blocks * num_threads; //为了数据分配的存储器大小，这里每一个人线程计算一个flaot
	//在host端分配内存，h_表示host端，i表示input，o表示output
	float* h_idata = nullptr;	float* h_odata = nullptr;
	h_idata =(float *)malloc(mem_size);
	h_odata = (float*)malloc(mem_size);	if(h_idata != nullptr) {	   memset(h_idata, 0, mem_size);
	}else{		return;
	}	if(h_odata!=nullptr){		memset(h_odata, 0, mem_size);
	}else{		return;
	}	
	//在device端分配显存，d_表示device端，i表示input，o表示output
	float* d_idata = nullptr;soccer.xj.zuq8.info	float* d_odata = nullptr;
	cudaError_t cudaStatus;  //cuda状态判断
	cudaStatus=cudaMalloc((void**)&d_idata, mem_size);	if(cudaStatus != cudaSuccess){		printf("d_idata is cudaMalloc failed!\n");		return;
	}
	cudaStatus=cudaMalloc((void**)&d_odata, mem_size);	if(cudaStatus!=cudaSuccess){		printf("d_odata is cudaMalloc failed!\n");		return;
	}	
	//初始化内存中的值
	for(unsigned int i = 0; i < num_threads * num_blocks;i++){
		h_idata[i] =1.0f;
	}//end for(unsigned int i = 0; i < num_threads * num_blocks;i++)
	//将内存中的输入数据读入设备端显存，这样就完成了主机对设备的数据写入
	cudaStatus=cudaMemcpy(d_idata,h_idata,mem_size,cudaMemcpyHostToDevice);	//设置运行参数，即网格的形状和线程块的形状
	dim3 grid(num_blocks,1,1);
	dim3 block(num_threads,1,1);	// 运行核函数，调用GPU进行运算
	testMatrixAssignKernel <<<grid, block>>> (d_idata,d_odata);	//将结果从显存写入内存
	cudaStatus = cudaMemcpy(h_odata,d_odata,mem_size,cudaMemcpyDeviceToHost);	//打印结果
	printf("赋值前的矩阵：\n");	for (unsigned int iblock = 0; iblock < num_blocks; iblock++) {		for (unsigned int ithread = 0; ithread < num_threads; ithread++) {			printf("%5.0f", h_idata[iblock * num_threads + ithread]);
		}//end for(unsigned int ithread = 0; ithread < num_threads; ithread++)
		printf("\n");
	}//end for(unsigned int iblock = 0; iblock < num_blocks; iblock++)
	printf("赋值后的矩阵：\n");	for(unsigned int iblock = 0; iblock < num_blocks; iblock++){		for(unsigned int ithread = 0; ithread < num_threads; ithread++){			printf("%5.0f",h_odata[iblock*num_threads+ithread]);
		}//end for(unsigned int ithread = 0; ithread < num_threads; ithread++)
		printf("\n");
	}//end for(unsigned int iblock =gov.js.zuzb.info 0; iblock < num_blocks; iblock++)
	//输出存储器指针
	free(h_idata);	free(h_odata);
	cudaFree(d_idata);
	cudaFree(d_odata);
}