一个.cu文件内既包含CPU程序（称为主机程序），也包含GPU程序（称为设备程序）。如何区分主机程序和设备程序？根据声明，凡是挂有“global”或者“device”前缀的函数，都是在GPU上运行的设备程序，不同的是__global__设备程序可被主机程序调用，而__device__设备程序则只能被设备程序调用。

CUDA程序文件后缀为.cu，有些编译器可能不认识这个后缀的文件，我们可以在VS的Tools->Options->Text Editor->File Extension里添加cu后缀到VC++中

CUDA 线程模型

CUDA的线程组织结构。首先我们都知道，线程是程序执行的最基本单元，CUDA的并行计算就是通过成千上万个线程的并行执行来实现的。下面的机构图说明了GPU的不同层次的结构。

CUDA的线程模型从小往大来总结就是：

Thread：线程，并行的基本单位
Thread Block：线程块，互相合作的线程组，线程块有如下几个特点：
允许彼此同步
可以通过共享内存快速交换数据
以1维、2维或3维组织
Grid：一组线程块
以1维、2维组织
共享全局内存
Kernel：在GPU上执行的核心程序，这个kernel函数是运行在某个Grid上的。

One kernel <-> One Grid
每一个block和每个thread都有自己的ID，我们通过相应的索引找到相应的线程和线程块。

threadIdx，blockIdx
Block ID: 1D or 2D
Thread ID: 1D, 2D or 3D
理解kernel，必须要对kernel的线程层次结构有一个清晰的认识。首先GPU上很多并行化的轻量级线程。kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。线程两层组织结构如上图所示，这是一个gird和block均为2-dim的线程组织。grid和block都是定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1。因此grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构，kernel调用时也必须通过执行配置<<<grid, block>>>来指定kernel所使用的网格维度和线程块维度。举个例子，我们以上图为例，分析怎么通过<<<grid,block>>>>这种标记方式索引到我们想要的那个线程。CUDA的这种<<<grid,block>>>其实就是一个多级索引的方法，第一级索引是(grid.xIdx, grid.yIdy)，对应上图例子就是(1, 1)，通过它我们就能找到了这个线程块的位置，然后我们启动二级索引(block.xIdx, block.yIdx, block.zIdx)来定位到指定的线程。这就是我们CUDA的线程组织结构。

SP和SM（流处理器）

SP：最基本的处理单元，streaming processor，也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。
SM：多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。
需要指出，每个SM包含的SP数量依据GPU架构而不同，Fermi架构GF100是32个，GF10X是48个，Kepler架构都是192个，Maxwell都是128个。

简而言之，SP是线程执行的硬件单位，SM中包含多个SP，一个GPU可以有多个SM（比如16个），最终一个GPU可能包含有上千个SP。这么多核心“同时运行”，速度可想而知，这个引号只是想表明实际上，软件逻辑上是所有SP是并行的，但是物理上并不是所有SP都能同时执行计算（比如我们只有8个SM却有1024个线程块需要调度处理），因为有些会处于挂起，就绪等其他状态，这有关GPU的线程调度。

下面这个图将从硬件角度和软件角度解释CUDA的线程模型。

每个线程由每个线程处理器（SP）执行
线程块由多核处理器（SM）执行
一个kernel其实由一个grid来执行，一个kernel一次只能在一个GPU上执行

block是软件概念，一个block只会由一个sm调度，程序员在开发时，通过设定block的属性，告诉GPU硬件，我有多少个线程，线程怎么组织。而具体怎么调度由sm的warps scheduler负责，block一旦被分配好SM，该block就会一直驻留在该SM中，直到执行结束。一个SM可以同时拥有多个blocks，但需要序列执行。下图显示了GPU内部的硬件架构：

CUDA内存模型

CUDA中的内存模型分为以下几个层次：

每个线程都用自己的registers（寄存器）
每个线程都有自己的local memory（局部内存）
每个线程块内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源
每个grid都有自己的global memory（全局内存），不同线程块的线程都可使用
每个grid都有自己的constant memory（常量内存）和texture memory（纹理内存），），不同线程块的线程都可使用
线程访问这几类存储器的速度是register > local memory >shared memory > global memory

下面这幅图表示就是这些内存在计算机架构中的所在层次。

CUDA编程模型

CUDA术语：

global

通过关键字就可以表示某个程序在CPU上跑还是在GPU上跑！如下表所示，比如我们用__global__定义一个kernel函数，就是CPU上调用，GPU上执行，注意__global__函数的返回值必须设置为void。

CPU和GPU间的数据传输

首先介绍在GPU内存分配回收内存的函数接口：

cudaMalloc(): 在设备端分配global memory
cudaFree(): 释放存储空间
CPU的数据和GPU端数据做数据传输的函数接口是一样的，他们通过传递的函数实参（枚举类型）来表示传输方向：

cudaMemcpy(void *dst, void *src, size_t nbytes,
enum cudaMemcpyKind direction)

enum cudaMemcpyKind:

cudaMemcpyHostToDevice（CPU到GPU）
cudaMemcpyDeviceToHost（GPU到CPU）
cudaMemcpyDeviceToDevice（GPU到GPU）

用代码表示线程组织模型

我们可以用dim3类来表示网格和线程块的组织方式，网格grid可以表示为一维和二维格式，线程块block可以表示为一维、二维和三维的数据格式。

dim3 DimGrid(100, 50);  //5000个线程块，维度是100*50
dim3 DimBlock(4, 8, 8);  //每个线层块内包含256个线程，线程块内的维度是4*8*8

怎么计算线程号

1.使用N个线程块，每一个线程块只有一个线程，即

dim3 dimGrid(N);
dim3 dimBlock(1);

此时的线程号的计算方式就是

threadId = blockIdx.x;

其中threadId的取值范围为0到N-1。对于这种情况，我们可以将其看作是一个列向量，列向量中的每一行对应一个线程块。列向量中每一行只有1个元素，对应一个线程。

2.使用M×N个线程块，每个线程块1个线程

dim3 dimGrid(M,N);
dim3 dimBlock(1);

blockIdx.x 取值0到M-1
blcokIdx.y 取值0到N-1

这种情况一般用于处理2维数据结构，比如2维图像。每一个像素用一个线程来处理，此时需要线程号来映射图像像素的对应位置，如

pos = blockIdx.y * blcokDim.x + blockIdx.x; //其中gridDim.x等于M

3.使用一个线程块，该线程具有N个线程，即

dim3 dimGrid(1);
dim3 dimBlock(N);

此时线程号的计算方式为
threadId = threadIdx.x;

其中threadId的范围是0到N-1，对于这种情况，可以看做是一个行向量，行向量中的每一个元素的每一个元素对应着一个线程。

4.使用M个线程块，每个线程块内含有N个线程，即

dim3 dimGrid(M);
dim3 dimBlock(N);

这种情况，可以把它想象成二维矩阵，矩阵的行与线程块对应，矩阵的列与线程编号对应，那线程号的计算方式为

threadId = threadIdx.x + blcokIdx*blockDim.x;

上面其实就是把二维的索引空间转换为一维索引空间的过程。

5.使用M×N的二维线程块，每一个线程块具有P×Q个线程，即

dim3 dimGrid(M, N);
dim3 dimBlock(P, Q);

这种情况其实是我们遇到的最多情况，特别适用于处理具有二维数据结构的算法，比如图像处理领域。

其索引有两个维度

threadId.x = blockIdx.x*blockDim.x+threadIdx.x;
threadId.y = blockIdx.y*blockDim.y+threadIdx.y;

上述公式就是把线程和线程块的索引映射为图像像素坐标的计算方法。

CUDA应用例子

https://www.cnblogs.com/mtcnn/p/9411877.html

https://www.cnblogs.com/skyfsm/p/9673960.html

REF
https://blog.csdn.net/qq_30263737/article/details/81235580

CUDA编程.cu文件相关推荐

matlab如何使用cu文件,Matlab编译cuda的.cu文件
matlab函数,大体首先是用nvcc命令生成.o文件,然后用mex链接对应库文件,生成动态链接库(.mexw64等). 测试环境: 1) Windows x64 + matlab + cuda 5. ...
matlab cuda的.cu文件应该放在那里_无人机基于Matlab/Simulink的模型开发(连载一）
"一切可以被控制的对象,都需要被数学量化" 这是笔者从事多年研发工作得出的道理,无论是车辆控制,机器人控制,飞机控制,还是无人机控制,所有和机械运动相关的控制,如果不能被很好的数学 ...
matlab cuda的.cu文件应该放在那里_App Store限免推荐｜日历和文件管理、无尽狂奔、Scalak等，共5款...
在今天限免的 iOS 应用中,我们精选了以下 3 款限免应用,2 款限免游戏.识别二维码/复制 logo 上方名字即可前往 AppStore 下载.如遇恢复原价,则表示限免已结束,请谨慎下载. 具体下 ...
.cu文件的两种编译方式
开一个cuda学习的坑,记录自己学习cuda过程中一些有意思的知识点主要参考教程:谭升的博客 CUDA文件既可以使用CMake编译,也可以使用nvcc编译.区别在于,CMake能使用单条命令进行多个 ...
深入浅出CUDA编程
标签: cuda编程threadfloatconflictexpress 2010-12-10 13:29 44960人阅读评论(7) 收藏举报 CUDA 是 NVIDIA 的 GPGPU 模型, ...
CUDA编程指南阅读笔记
随着多核CPU和众核GPU的到来,并行编程已经得到了业界越来越多的重视,CPU-GPU异构程序能够极大提高现有计算机系统的运算性能,对于科学计算等运算密集型程序有着非常重要的意义.这一系列文章是根据& ...
CUDA 编程上手指南：CUDA C 编程及 GPU 基本知识
作者丨科技猛兽编辑丨极市平台本文原创首发于极市平台,转载请获得授权并标明出处. 推荐大家关注极市平台公众号,每天都会更新最新的计算机视觉论文解读.综述盘点.调参攻略.面试经验等干货~ 目录 1 C ...
CUDA 编程实例:计算点云法线
程序参考文章:http://blog.csdn.net/gamesdev/article/details/17535755 程序优化2 简介:CUDA ,MPI,Hadoop都是并行运算的工具.CU ...
【转载】cuda编程入门
目录 1.什么是CUDA 2.为什么要用到CUDA 3.CUDA环境搭建 4.第一个CUDA程序 5. CUDA编程 5.1. 基本概念 5.2. 线程层次结构 5.3. 存储器层次结构 5.4. 运 ...
CUDA编程：与OpenCV结合
前言学习计算机图像处理算法的童鞋,就不得不学习cuda,为啥呢?因为图像处理一般都是矩阵运算,动不动就是百万的计算量这个时候优化计算时间是必不可少的.openCV本身提供了很多cuda函数,能够满足 ...

CUDA编程.cu文件