赞助商链接
本类月下载排行
本类周下载排行
常用软件推荐
CUDA调试工具(SDK编译器)可以对所有NVIDIA的GPU进行控制并实时获得GPU状态和运行情况。
CUDA编程笔记
1、下载CUDA工具包、驱动和SDK，依次安装，编译SDK里的工程文件，具体配置方法可谷歌百度，基本步骤就是先编译sdk里面的几个库，然后设置代码相应库函数关键字高亮。
2、 下载CUDA调试工具Nsight，可以实现单卡双机调试或双卡单机调试，虽然听说2.2可以单卡单机，但还没找到方法。具体调试可参考Nsight的help文档和上的别人笔记（CUDA双机Nsight调试总结）。
3、 Nsight2.2中包含cuda4.0，cuda4.1， cuda4.2 rules。
4、如果在运行cuda程序时出现黑屏花屏显卡驱动奔溃几秒然后重新启动，个人认为应该是显存不足造成的,后来发现不是这个原因，这时候打开Nsight 的Monitor，将Options/General下的WDDM TDR enabled 设置为 False，重启电脑，这时再运行程序，显示器不会更新画面，GPU会提供完全的显存专注于运行我们的程序。直到运行结束，显示器画面恢复更新。这种方法当然只有用在实验上，不能用在实际中。
5、 普通C++控制台和MFC程序中如何新建调用CUDA程序（VS2010）：
1） 首先新建一个C++工程，然后右击工程，打开&生成自定义&，选择最新的cudarules。
2） 然后新建两个.cu文件，分别命名为kernel.cu和func.cu，前者存放核函数，后者存放调用核函数的基本cuda函数。然后打开两个文件的属性，将项类型设置为CUDA C/C++，然后在CUDAC/C++下的Device里的Code Generation参数修改为运行电脑GPU的计算等级，例如1.1的就写成compute_11,sm_11。还有，前者kernel.cu属性中&从生成中排除&要选&是&。
3） 在func.cu文件中将kernel.cu当做头文件包含，里面的函数进行基本的cuda程序操作：分配显卡内存、拷贝内存到显存、调用kernel函数、拷贝显存到内存。
4） 在func.cu文件中定义函数时，将func函数定义为C扩展函数：extern &C&void func();
5） 新建一个cpp文件，作为主程序，先在程序上面添加C扩展函数func（）的说明，此时不需要include &func.cu&。然后再下面的代码中调用func函数即可。
6） 打开整个工程的属性，在CUDA C/C++下的Device里的Code Generation参数依然修改为运行电脑GPU的计算等级。
7） 在工程属性的&链接器/附加依赖项&中添加cudart.
6、 CUDA Tookit中的NVIDIAVisual Profiler是个不错的cuda程序性能分析软件，能生成程序中GPU处理部分的各个步骤的时间条，以及各个阶段的具体参数。
7、 CUDA提供了CUBLA（具有基本的向量和矩阵运算）库，SDK中已经自带了，不用另外下载。它的运算函数分为三个级别：1）向量与向量的运行2）矩阵与向量的运算3）矩阵与矩阵的运算。
工程属性的链接器中的附加依赖项需添加：cublas.lib。
8、 在CUBLAS中矩阵是按列优先存储的，这与C中按行存储矩阵的思想正好相反。
基本的调用步骤举例：
1） 包含cublas的头文件：
2） 创建一个cublas上下文句柄：
3） 调用cublas函数处理向量或矩阵运算：
4） 销毁cublas句柄：
关於CUDA在VS2008下的调试问题
在前提只有一台主机一张NV显卡的情况下
不知道大家都怎N调试自己的程序...
在VS总是找不到一个有效率的程序去调试他...
用VS就是榱撕玫魇
不过...NV自己出的那套Nsight貌似要两张显卡或者两台主机才能调试
&&请点击以下链接下载该软件：&CUDA调试工具(SDK编译器)V2.0.1.2 绿色版
上一软件：
下一软件：
(评论内容只代表网友观点，与本站立场无关)
为了保证您快速的下载,推荐使用[] 、[] 等专业工具下载.
为确保软件能正常使用,请使用[
]解压本站软件.
目前不少软件都捆绑流氓插件，请在安装的之时务必留意每一个安装步骤.绿色先锋本身是不会捆绑任何插件在软件中的.
该软件为网上收集，若无意中侵犯了您的版权，.我们将在收信后24小时内删除侵权内
本站下载的软件中，部分软件经过压缩加密处理，解压密码为：
感谢您对绿色先锋的支持，请将网站地址放在您的博客，空间等地方，以便我们为您及您的朋友提供更好的服务.
软件按字母排列：
中文按声母搜索：只需一步，快速开始
查看: 160|回复: 1
cudaer 分享并行cuda代码调试心得(对于不使用IDE的情况下)特指在Linux非图形化界面
cuda 并行代码其实跟C语言差不多，主要是在于调试的过程中，出现很大的问题，由于很多的朋友没有很系统的学习有关nsight ，以及对cuda-gdb不是很熟的情况下，我们将采取下面的办法来调试代码，这样做的效率可能不高，但是对于新手，还是建议手动找错，有利于熟悉整个cuda代码运行的流程，同时也能对整个代码的流程很清楚。具体的流程有下面几部分（特指在Linux非图形化界面的情况下）：
（1）使用 nvcc 来编译代码，如果发现一些语法错误，先解决一些语法错误，从而使最终的代码没有出现语法错误为止，这里需要相当注意。然后我们再说明一些逻辑错误的问题。
（2）cuda 代码中，我们一般使用cuda_runtime api 和cuda_driver api 对于新手，我们一般学会使用cuda_runtime&&api 就可以了，对于cuda_runtime 的一些函数（一般都是以cuda*开头的函数），我们在使用的时候，需要检查一下，每次调用的返回值，来确定是不是调用成功了，这里有两种方法， 第一种使用简单麻烦的语句 ： 比如 cudaError_t&& error=cuda*(); 然后通过判断error的值是否是cudaSl来确定函数调用是否成功了。第二种方法是使用宏定义的方法，定义一个判断语句是否成功的语句，然后，我们可以使用下面的宏定义来确认每一次调用函数，来判断是否成功 代码如下：
#define CUDA_CALL(cuda_function, ...)&&{ \
& & cudaError_t status = cuda_function(__VA_ARGS__); \
& & cudaEnsureSuccess(status, #cuda_function, false, __FILE__, __LINE__); \
bool cudaEnsureSuccess(cudaError_t status, const char* status_context_description,
& && &&&bool die_on_error, const char* filename, unsigned line_number) {
& & if (status_context_description == NULL)
& && &&&status_context_description = &&;
& & if (status == cudaSuccess) {
#if REPORT_CUDA_SUCCESS
& && && &cerr &&&&&Succeeded: & && status_context_description && std::endl && std::
& && &&&return
& & const char* errorString = cudaGetErrorString(status);
& & cerr && &CUDA Error: &;
& & if (status_context_description != NULL) {
& && &&&cerr && status_context_description && &: &;
& & if (errorString != NULL) {
& && &&&cerr && errorS
& & else {
& && &&&cerr && &(Unknown CUDA status code & && status && &)&;
& & if (filename != NULL) {
& && &&&cerr && filename && &:& && line_
& & cerr && std::
& & if(die_on_error) {
& && &&&exit(EXIT_FAILURE);
& && && && &// ... or cerr && &FATAL ERROR& && etc. etc.
在使用的时候，我们仅仅通过 CUDA_CALL(cuda*())来进行判断，在这里，就可以检查出这个cuda_runtime API 是不是调用成功了
（3）在检查核函数的调用的时候，我们需要进行两个方面的检查&&第一个方面，我们需要检查有关核函数运行完毕的时候就进行相关检查 代码如下：
kernel&&&&&&();
cudaError_t&&error=cudaGetLastError();
printf(&cuda Error: %s&,cudaLastErrorString(error));//这个地方主要是判断核函数运行后会残生什么样的错误。
//在进行数值传递的时候，特别在使用cudaMemcpy() 的时候，我们也还是需要检查这个cuda*（ ）运行时候的状态。
error=cudaMemcpy();
printf(&cuda Error: %s&,cudaLastErrorString(error)); // 这个地方是第二次判断在传递的时候，时候会产生错误。
从而确认时候会出现这个核函数的使用错误，在核函数的使用的过程中，我们往往会出现传递值的时候，会出现要计算的值不是一一对应的情况，同时也会出现越界的情况。
（4）在我们将我们需要使用的核函数进行相关的检查的时候，我们 往往会出现，一些错误会在你最后调用核函数的时候，出现报错的额情况，但是我们是真的以为是最后，那个核函数出现错误吗？ 答案是否定的，cuda 核函数运行的时候，往往会讲前面的错误进行相关传递，从而导致，我们及时在后面出现的错误，其实是前面的一些核函数产生的结果，这个时候，我们需要使用分割的调试的方法，即把核函数与核函数之间给隔离出来，然后独自的进行相关分析，我们就可以确定相关出现问题的地方确定在哪个地方，然后我们在修改这个地方，在这里，我们使用的方法如下面所示：
cudaGetLastError()&&//主要是用来获取上面代码运行的错误，从而，隔开错误，避免对后面的核函数，进行影响，这点我们是需要相当的重视。
cudaDeviceReset()
printf(&cat\n&);
kernel2&&&&&&();
printf(&dog/n&);
然后我们就通过kernel函数在进行调用时候，然后看看输出的结果有没有相关的改变，如果，最后的错误的输出没有报错，说明，这块是没有问题的。
那么错误可能会在前面。
然后重复这个动作，我们可以将错误确定到一个或几个函数的调用。
（5）在我们发现一个或者几个函数的调用失败的时候，我们需要先系统看看数据传递是否是正常的，如果数据传递是正常的，我们可以通过一些简单的数据来分析每部分运行的结果，我们可以确认到底是那部分函数有语法错误，然后再进行修改，一般在使用核函数，出现数据传递的问题和越界的现象很严重，要考虑自己的block 和grid 分块，对于自己的要进行数据会不会出现没有完全处理（即处理不完全）或者也是否出现在进行计算的时候（是否出现越界的情况，即在进行数据处理的时候获取一些不存在的数据，所以造成问题），然后分析，对于越界的情况我们需要添加一些判断的条件，从而避免出现越界，在出现对数据处理不足的情况，我们需要，将那些没有进行处理的数据，添加进去，进行计算。最后进行测试。
最后我总结一波：
我所分享的是由猫叔教导的，我通过实际应用，确实对我代码调试BUG存在很大的作用，在这感谢LADY姐 猫叔。
大家如果有关于cuda 并行计算的问题 欢迎加入
& &GPU应用开发技术群
楼主这文章有几个问题，
一是否少了一处#endif?
二是否根据你的宏，应当写成CUDA_CALL(cuda开头的函数名, 参数1, 参数2...)，而不是CUDA_CALL(cuda开头的函数(...))?
三是你在&&&&&&后面立刻跟随的cudaGetLastError是检测是否启动成功（例如无效的形状或者shared memory配置可能会导致启动失败），而不是运行中的错误。而后续的cudaMemcpy则因为自带的同步效果，可以返回的是运行中的问题。
四是有一处错别字：残生。这应当是产生。
等等吧。修正下？
站长推荐 /1
为了更好地为各位会员提供技术支持服务，本论坛采用注册审核制度。您注册会员后，我们的管理人员会在8小时内在后台进行审核，审核通过会发邮件通知。只要是真实的开发者或学习者，我们都会审核通过。敬请谅解！如果您着急的话，可以发邮件至： 提醒我们第一时间审核通过。博客访问： 1396658
博文数量： 367
博客积分： 4817
博客等级： 上校
技术积分： 4216
注册时间：
IT168企业级官微
微信号：IT168qiye
系统架构师大会
微信号：SACC2013
分类： 高性能计算
初学CUDA，笔记错误之处在所难免，还请发现问题的诸位读者不吝赐教。
CUDA全称是Compute Unified Device Architecture，中文名称即统一计算设备架构，它是NVIDIA公司提出了一种通用的并行计算平台和编程模型。使用CUDA，我们可以开发出同时在CPU和GPU上运行的通用计算程序，更加高效地利用现有硬件进行计算。为了简化并行计算学习，CUDA为程序员提供了一个类C语言的开发环境以及一些其它的如FORTRAN、DirectCOmpute、OpenACC的高级语言/编程接口来开发CUDA程序。
我们知道，不同的GPU拥有不同的核心数目，在核心较多的系统上CUDA程序运行的时间较短，而在核心较少的系统上CUDA程序的执行时间较多。那么，CUDA是如何做到的呢？
并行编程的中心思想是分而治之：将大问题划分为一些小问题，再把这些小问题交给相应的处理单元并行地进行处理。在CUDA中，这一思想便体现在它的具有两个层次的问题划分模型。一个问题可以首先被粗粒度地划分为若干较小的子问题，CUDA使用被称为块（Block）的单元来处理它们，每个块都由一些CUDA线程组成，线程是CUDA中最小的处理单元，将这些较小的子问题进一步划分为若干更小的细粒度的问题，我们便可以使用线程来解决这些问题了。对于一个普通的NVIDIA GPU，其CUDA线程数目通常能达到数千个甚至更多，因此，这样的问题划分模型便可以成倍地提升计算机的运算性能。
GPU是由多个流水多处理器构成的，流水处理器以块（Block）为基本调度单元，因此，对于流水处理器较多的GPU，它一次可以处理的块（Block）更多，从而运算速度更快，时间更短。而反之对于流水处理器较少的GPU，其运算速度便会较慢。这一原理可以通过下图形象地看出来：
本节将介绍CUDA的一些基本的编程概念，该节用到的例子来自于CUDA Sample中的VectorAdd项目。
CUDA C是C语言的一个扩展，它允许程序员定义一种被称为内核函数（Kernel Functions）的C函数，内核函数运行在GPU上，一旦启动，CUDA中的每一个线程都将会同时并行地执行内核函数中的代码。
内核函数使用关键字__global__来声明，运行该函数的CUDA线程数则通过<<>>执行配置语法来设置。（参见章节"C语言扩展"），每一个执行内核函数的线程都由一个唯一的线程ID，这一ID可以通过在内核函数中访问threadIdx变量来得到。
下面通过一些示例代码来展示刚刚提到的这些概念该如何应用在编程中：
在上面的代码中，N个线程将会并行地同时执行加法运算。
CUDA的每一个线程都有其线程ID，线程的ID信息由变量threadIdx给出。threadIdx是CUDA C语言的内建变量，通常它用一个三维数组来表示。使用三维数组的方便之处在于可以很方便地表示一维、二维和三维线程索引，进而方便地表示一维、二维和三维线程块（thread block）。这样，无论是数组、矩阵还是体积的计算，都可以很容易地使用CUDA进行运算。
线程的索引与线程ID之间存在着直接的换算关系，对于一个索引为(x, y, z)的线程来说：
每个线程块（block）中的线程数量是有限制的，因为依据前面所说，同一线程块（block）中的所有线程都会被分配到同一个处理器核上运行，共享有限的存储资源，因此对于当前的GPU，线程块所能包含的最大线程数目为1024。
上面的例子中numBlocks代表线程块的数量，这里的值为1。在一般的CUDA程序中，这个值通常大于1，也就是说将会有多个线程块被分配到多个处理器核中同时进行处理，这样就大大提高了程序的并行性。
在CUDA中，线程块包含在线程格（grid）当中，线程格可以是一维、二维或者三维的，线程格的尺寸一般根据待处理数据的规模或者处理器的数量来指定。线程格中所包含的线程块数目通常远远大于GPU处理器核心的数目。下图展示了线程格（grid）、线程块（block）以及线程（thread）之间的关系：
内核函数的调用可以简化为kernel<<>>(parameters)，在尖括号中，A代表线程格（grid）的尺寸，它可以是三维的，用类型dim3表示，也可以是一维的，用int类型表示。B代表线程块（block）的尺寸，它与A类似，也可分别用dim3或int类型表示。
在内核函数内部，CUDA为我们内建了一些变量用于访问线程格、线程块的尺寸和索引等信息，它们是：
& & & 1. gridDim：代表线程格（grid）的尺寸，gridDim.x为x轴尺寸，gridDim.y、gridDim.z类似。拿上图来说，它的gridDim.x = 3，gridDim.y = 2，gridDim.z = 1。
& & & 2. blockIdx：代表线程块（block）在线程格（grid）中的索引值，拿上图来说，Block（1,1）的索引值为：blockIdx.x = 1，blockIdx.y = 1。
& & & 3. blockDim：代表线程块（block）的尺寸，blockDIm.x为x轴尺寸，其它依此类推。拿上图来说，注意到Block（1,1）包含了4 * 3个线程，因此blockDim.x = 4, blockDim.y = 3。
& & & 4. threadIdx：线程索引，前面章节已经详细探讨过了，这里不再赘述。
明白了这些变量的含义，那么下面的矩阵加法程序便不难理解了：
在上面的程序中，线程块（block）的尺寸是16x16，这是CUDA编程中一个非常普遍的选择。线程格（grid）包含了足够多的线程块（block）来进行计算。
线程块（block）是独立执行的，在执行的过程中线程块之间互不干扰，因此它们的执行顺序是随机的。
同一线程块中的线程可以通过访问共享内存（shared memory）或者通过同步函数__syncthreads()来协调合作。这些概念将在以后的章节中详细解释。
在GPU上CUDA线程可以访问到的存储资源有很多，每个CUDA线程拥有独立的本地内存（local Memory）；每一个线程块（block）都有其独立的共享内存（shared memory），共享内存对于线程块中的每个线程都是可见的，它与线程块具有相同的生存时间；同时，还有一片称为全局内存（global memory）的区域对所有的CUDA线程都是可访问的。
除了上述三种存储资源以外，CUDA还提供了两种只读内存空间：常量内存（constant memory）和纹理内存（texture memory），同全局内存类似，所有的CUDA线程都可以访问它们。对于一些特殊格式的数据，纹理内存提供多种寻址模式以及数据过滤方法来操作内存。这两类存储资源主要用于一些特殊的内存使用场合。
一个程序启动内核函数以后，全局内存、常量内存以及纹理内存将会一直存在直到该程序结束。下面是CUDA的内存层次图：
3.4 异构编程（Heterogeneous Programming）
CUDA的异构编程模型假定CUDA线程都运行在一个可被看做CPU协处理器的芯片上，这就使得CUDA内核函数可以和CPU端C程序的运行并行运行，从而加快程序的运行效率。为了达到这个效果，CUDA程序需要管理两大块由DRAM构成的内存区域：CPU端可以访问到的主机内存（host memory）以及GPU端供CUDA内核访问到的设备内存（device memory），设备内存主要由全局内存、常量内存以及纹理内存构成。现在，CUDA程序的运行机制便很明了了：CPU端代码生成原始数据，通过CUDA运行时函数库将这些原始数据传输到GPU上，在CPU端启动CUDA内核函数进行运算，然后将运算结果从设备端传输到主机端，计算任务便完成了。
4.语言编程接口
异构程序设计跟传统的串行程序设计差别是很大的，学习起来也是非常不容易的。NVIDIA非常够意思，为了简化CUDA的学习曲线，它采用了绝大多数程序员都熟悉的C语言作为其根基，CUDA C是NVIDIA为程序员提供的一类编程接口，它实际上是一个C语言的扩展，在C的基础上增加了一些新的语法和变量，并且提供了功能丰富的库函数，方便程序员使用GPU进行异构计算。
除了前面章节提到的CUDA最基本、最核心的概念以外，CUDA C呈现给程序员的接口主要由两大类API构成，它们分别是CUDA Runtime API和CUDA Driver API，Runtime API实际上是对于Driver API的封装，其目的自然是方便程序员的代码编写工作。Driver API为用户提供了更细一层的控制手段，通过它可以控制诸如CUDA Contexts（一种类似主机进程的概念）以及CUDA Modules（类似主机动态加载库的概念）等更加底层的CUDA模块。
任何一种程序设计语言都需要相应的编译器将其编译为二进制代码，进而在目标机器上得到执行。对于异构计算而言，这一过程与传统程序设计语言是有一些区别的。为什么？因为CUDA它本质上不是一种语言，而是一种异构计算的编程模型，使用CUDA C写出的代码需要在两种体系结构完全不同的设备上执行：1、CPU；2、GPU。因此，CUDA C的编译器所做的工作就有点略多了。一方面，它需要将源代码中运行在GPU端的代码编译得到能在CUDA设备上运行的二进制程序。另一方面，它也需要将源代码中运行在CPU端的程序编译得到能在主机CPU上运行的二进制程序。最后，它需要把这两部分有机地结合起来，使得两部分代码能够协调运行。
CUDA C为我们提供了这样的编译器，它便是NVCC。严格意义上来讲，NVCC并不能称作编译器，NVIDIA称其为编译器驱动（Compiler Driver），本节我们暂且使用编译器来描述NVCC。使用nvcc命令行工具我们可以简化CUDA程序的编译过程，NVCC编译器的工作过程主要可以划分为两个阶段：离线编译（Offline Compilation）和即时编译（Just-in-Time Compilation）。
& & 下面这幅图简单说明了离线编译的过程：
CUDA C Runtime函数库没有明确的初始化函数，在程序第一次调用Runtime库函数时它会自动初始化。因此，在记录Runtime函数调用时间和理解程序中第一个Runtime调用返回的错误代码时，需要将初始化考虑在内。
片段展示了设备内存的分配，传输以及回收过程。
更多详细的内容请查阅参考手册。
使用cudaGetSymbolAddress()函数可以获得被声明存储在全局内存中的变量地址。为了获得分配内存的大小，可以使用cudaGetSymbolSize()函数。
可以看出，为了计算矩阵C的任何一个元素，程序都需要从全局内存（global memory）中获得矩阵A的一行和矩阵B的一列。因此，完成这一计算矩阵A被读取了B.width次，矩阵B被读取了A.height次。
& & & & 现在我们来使用共享内存（shared memory）实现矩阵乘法。假设矩阵C可以被划分为若干个较小的子方阵Csub，我们使用一个线程块（thread block）来负责某一子方阵的计算，线程块中的每一个线程（thread）正好负责子方阵Csub中一个元素的计算。这样划分后，任何一个结果子方阵Csub'（尺寸为block_size * block_size）都是与该方阵具有相同行索引的尺寸为A.width * block_size的A的子矩阵Asub和与该方阵具有相同列索引的尺寸为block_size * B.height的B的子矩阵Bsub相乘所得到。
和Bsub被划分为尽可能多的分离的维度为block_size的子方阵，Csub的值便是这些子矩阵相乘后相加所得到的结果。子矩阵乘法的执行顺序都是首先将它们从全局内存（global memory）拷贝到共享内存（shared memory）（线程块中的每一个线程正好负责方阵一个元素的拷贝），然后由线程自己完成相应元素的计算任务，利用寄存器存储局部结果，最后将寄存器的内容与新得到的计算结果依此累加起来得到最终运算结果并将其传输到全局内存（global memory）中。
下面的代码定义了每一个流的行为：从主机端拷贝数据到设备端，内核启动，从设备端拷贝数据到主机端：
这部分代码中有一点需要注意：为了并行化数据拷贝和内核执行，主机端内存必须分配为锁页（page-locked）内存。
要销毁一个流需要调用函数
cudaStreamDestroy()函数等待之前流中的指令序列运行完成，然后销毁指定流，将控制权返还给主机端。
可将stream[0]的内核启动和stream[1]从主机端到设备端的数据拷贝重叠起来并行执行。
上面的代码定义了两个流的操作，每个流都完成一次主机端到设备端的数据拷贝，一次内核启动，一次设备端到主机端的数据拷贝，最后增加了一个加入回调函数的操作。当设备端代码运行到回调函数点的时候，设备将控制权交还给主机端，主机端运行完成以后再将控制权返还给设备端，然后设备端继续运行。
阅读(12363) | 评论(0) | 转发(1) |
相关热门文章
给主人留下些什么吧！~~
请登录后评论。[菜鸟每天来段CUDA_C]使用多个CUDA流提高程序执行效率 - 推酷
[菜鸟每天来段CUDA_C]使用多个CUDA流提高程序执行效率
CUDA流表示一个GPU操作队列，并且该队列中的操作以添加到队列的先后顺序执行。使用CUDA流可以实现任务级的并行，比如当GPU在执行核函数的同时，还可以在主机和设备之间交换数据(前提是GPU支持重叠，property的deviceOverlay为true)。
cudaMemcpyAsync函数的功能是在GPU和主机之间复制数据。它是一个异步函数，即函数被调用后，只是放置一个请求，表示在流中执行一次内存复制操作。函数返回时，复制操作不一定启动或执行结束，只是该操作被放入执行队列，在下一个被放入流中的操作之前执行。
实验通过把一组数据分块复制到GPU执行，返回执行结果，来说明使用cuda流的使用能提高程序的执行效率。原理主要是使数据复制操作和核函数执行操作交叉执行，不用等到第一次核函数执行结束再开始第二轮的数据复制，以减少顺序执行带来的延迟(类似于编译中使用流水线在解决冲突的前提下提高效率)。
程序代码如下：
#include &cuda_runtime.h&
#include &cutil_inline.h&
#include &stdio.h&
#include &math.h&
static void HandleError( cudaError_t err,const char *file,int line )
if (err != cudaSuccess)
printf( &%s in %s at line %d\n&, cudaGetErrorString( err ),
file, line );
exit( EXIT_FAILURE );
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
#define N ()
#define FULL_DATA_SIZE N*20
__global__ void kernel(int* a, int *b, int*c)
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int offset = gridDim.x * blockDim.x;
if (idx & N)
int idx1 = (idx + 1) % 256;
int idx2 = (idx + 2) % 256;
float as = (a[idx] + a[idx1] + a[idx2]) / 3;
float bs = (b[idx] + b[idx1] + b[idx2]) / 3;
c[idx] = (as + bs) / 2;
int main()
cudaDeviceP
int devID;
HANDLE_ERROR(cudaGetDevice(&devID));
HANDLE_ERROR(cudaGetDeviceProperties(&prop, devID));
if (!prop.deviceOverlap)
printf(&No device will handle overlaps. so no speed up from stream.\n&);
cudaEvent_t start,
float elapsedT
HANDLE_ERROR(cudaEventCreate(&start));
HANDLE_ERROR(cudaEventCreate(&stop));
HANDLE_ERROR(cudaEventRecord(start, 0));
cudaStream_t stream0;
cudaStream_t stream1;
HANDLE_ERROR(cudaStreamCreate(&stream0));
HANDLE_ERROR(cudaStreamCreate(&stream1));
int *host_a, *host_b, *host_c;
int *dev_a0, *dev_b0, *dev_c0;
int *dev_a1, *dev_b1, *dev_c1;
HANDLE_ERROR(cudaMalloc((void**)&dev_a0, N*sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**)&dev_b0, N*sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**)&dev_c0, N*sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**)&dev_a1, N*sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**)&dev_b1, N*sizeof(int)));
HANDLE_ERROR(cudaMalloc((void**)&dev_c1, N*sizeof(int)));
HANDLE_ERROR(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault));
HANDLE_ERROR(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault));
HANDLE_ERROR(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault));
for (int i=0; i&FULL_DATA_SIZE; i++)
host_a[i] = rand();
host_b[i] = rand();
// tasks are put into stack for gpu execution
for (int i=0; i&FULL_DATA_SIZE; i+=2*N)
HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a+i, N*sizeof(int), cudaMemcpyHostToDevice, stream0));
HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1));
HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b+i, N*sizeof(int), cudaMemcpyHostToDevice, stream0));
HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1));
kernel&&&N/256, 256, 0, stream0&&&(dev_a0, dev_b0, dev_c0);
kernel&&&N/256, 256, 0, stream1&&&(dev_a1, dev_b1, dev_c1);
HANDLE_ERROR(cudaMemcpyAsync(host_c+i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0));
HANDLE_ERROR(cudaMemcpyAsync(host_c+i+N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1));
HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a+i, N*sizeof(int), cudaMemcpyHostToDevice, stream0));
HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b+i, N*sizeof(int), cudaMemcpyHostToDevice, stream0));
kernel&&&N/256, 256, 0, stream0&&&(dev_a0, dev_b0, dev_c0);
HANDLE_ERROR(cudaMemcpyAsync(host_c+i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0));
HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1));
HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1));
kernel&&&N/256, 256, 0, stream1&&&(dev_a1, dev_b1, dev_c1);
HANDLE_ERROR(cudaMemcpyAsync(host_c+i+N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1));
// wait until gpu execution finish
HANDLE_ERROR(cudaStreamSynchronize(stream0));
HANDLE_ERROR(cudaStreamSynchronize(stream1));
HANDLE_ERROR(cudaEventRecord(stop, 0));
HANDLE_ERROR(cudaEventSynchronize(stop));
HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));
printf(&Time taken: %3.1f ms\n&, elapsedTime);
// free stream and mem
HANDLE_ERROR(cudaFreeHost(host_a));
HANDLE_ERROR(cudaFreeHost(host_b));
HANDLE_ERROR(cudaFreeHost(host_c));
HANDLE_ERROR(cudaFree(dev_a0));
HANDLE_ERROR(cudaFree(dev_b0));
HANDLE_ERROR(cudaFree(dev_c0));
HANDLE_ERROR(cudaFree(dev_a1));
HANDLE_ERROR(cudaFree(dev_b1));
HANDLE_ERROR(cudaFree(dev_c1));
HANDLE_ERROR(cudaStreamDestroy(stream0));
HANDLE_ERROR(cudaStreamDestroy(stream1));
相对于顺序执行，使用两个cuda流使程序的执行时间少了20ms（由于数据量不大，所以使用流的优势不太明显）.
已发表评论数()
请填写推刊名
描述不能大于100个字符!
权限设置： 公开
仅自己可见
正文不准确
标题不准确
排版有问题
主题不准确
没有分页内容
图片无法显示
视频无法显示
与原文不一致}