页锁定内存:
cudaHostAlloc()分配页锁定内存,页锁定内存也称为固定内存或不可分页内存,它有一个重要的属性:操作系统将不会对这块内存分页并交换到磁盘上,从而确保了该内存始终驻留在物理内存中。
流:
cuda流用于任务的并行。任务并行性是指并行执行两个或多个不同的任务,而不是在大量数据上执行同一个任务的数据并行性。比如处理同一副图,你用一个流处理左边半张图片,再用第二个流处理右边半张图片,这两个流中的代码同时执行,加快了处理速度。
示例:
#include <stdio.h> #include <cuda_runtime.h> #define N (1024*1024) #define DATA_SIZE (N*20) __global__ void add(int *a,int *b,int *c){ int idx=threadIdx.x+blockIdx.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.0f; float bs=(b[idx]+b[idx1]+b[idx2])/3.0f; c[idx]=(as+bs)/2; } } int main(){ cudaDeviceProp prop; int whichDevice; cudaGetDevice(&whichDevice); cudaGetDeviceProperties(&prop,whichDevice); if(!prop.deviceOverlap){ printf("Device not overlap...."); return 0; } int *a,*b,*c; int *a1,*b1,*c1; int *a_host,*b_host,*c_host; cudaEvent_t start,end; float elapsedTime; cudaEventCreate(&start); cudaEventCreate(&end); cudaEventRecord(start,0); cudaStream_t stream0,stream1; cudaStreamCreate(&stream0); cudaStreamCreate(&stream1); cudaMalloc((void **)&a,N*sizeof(int)); cudaMalloc((void **)&b,N*sizeof(int)); cudaMalloc((void **)&c,N*sizeof(int)); cudaMalloc((void **)&a1,N*sizeof(int)); cudaMalloc((void **)&b1,N*sizeof(int)); cudaMalloc((void **)&c1,N*sizeof(int)); cudaHostAlloc((void **)&a_host,DATA_SIZE*sizeof(int),cudaHostAllocDefault); cudaHostAlloc((void **)&b_host,DATA_SIZE*sizeof(int),cudaHostAllocDefault); cudaHostAlloc((void **)&c_host,DATA_SIZE*sizeof(int),cudaHostAllocDefault); for(int i=0;i<DATA_SIZE;i++){ a_host[i]=i; b_host[i]=i; } for(int i=0;i<DATA_SIZE;i+=N*2){ cudaMemcpyAsync(a,a_host+i,N*sizeof(int),cudaMemcpyHostToDevice,stream0); cudaMemcpyAsync(a1,a_host+i+N,N*sizeof(int),cudaMemcpyHostToDevice,stream1); cudaMemcpyAsync(b,b_host+i,N*sizeof(int),cudaMemcpyHostToDevice,stream0); cudaMemcpyAsync(b1,b_host+i+N,N*sizeof(int),cudaMemcpyHostToDevice,stream1); add<<<N/256,256,0,stream0>>>(a,b,c); add<<<N/256,256,0,stream1>>>(a1,b1,c1); cudaMemcpyAsync(c_host+i,c,N*sizeof(int),cudaMemcpyDeviceToHost,stream0); cudaMemcpyAsync(c_host+i+N,c,N*sizeof(int),cudaMemcpyDeviceToHost,stream1); } cudaStreamSynchronize(stream0); cudaStreamSynchronize(stream1); cudaEventRecord(end,0); cudaEventSynchronize(end); cudaEventElapsedTime(&elapsedTime,start,end); printf("tie===%3.1f ms\n",elapsedTime); cudaFreeHost(a_host); cudaFreeHost(b_host); cudaFreeHost(c_host); cudaFree(a); cudaFree(b); cudaFree(c); cudaFree(a1); cudaFree(b1); cudaFree(c1); cudaStreamDestroy(stream0); cudaStreamDestroy(stream1); return 0; }
注:硬件在处理内存复制和核函数执行时分别采用了不同的引擎,因此我们需要知道,将操作放入流中队列中的顺序将影响着CUDA驱动程序调度这些操作以及执行的方式。 因此在将操作放入流的队列时应该采用宽度优先方式,而非深度优先方式。
参考:《GPU高性能编程CUDA实战》
相关推荐
例如,NVIDIA 的 Kepler 架构就具有 15 个流多处理器(SM),每个 SM 都有 6 个 issue unit 和 32 个 way SIMD。 Latency Optimized CPU CPU 的架构则是为了追求低延迟而设计的。CPU 的处理单元都是为串行任务而...
流 大事记 符号 使用纹理对象API的纹理 要求 CUDA≥7.0支持 C ++ 11编译器 CMake≥2.8.12 安装 cuda-wrapper使用CMake,标头仅安装到CMake指定的目录中: $ cmake path/to/cuda-wrapper $ make $ make
使用cuda-api-wrappers,您仍然可以拥有设备,流,事件等-但它们将以更多的C ++惯用方式更加方便地使用。主要特征所有函数和方法都会在失败时引发异常-无需检查返回值(异常带有状态信息)。 明智的命名空间(以及...
CUDA样本CUDA开发人员示例,演示了CUDA Toolkit中的功能。 该版本支持 。发行说明本部分仅描述GitHub上CUDA... 演示使用cudaMallocAsync和cudaMemPool系列API在GPU上按流排序的内存分配。 添加了streamOrderedAllocati
18645 学期项目(如何编写快速代码) @Shitao Weng : @Shushan Chen : 先决条件 ...标准输出流: PreEmp: 1.0210488 s , 1.36219% Windowing: 0.0832372 s , 5.38676% FFT padding: 0.018985 s ,
流实现不需要训练数据随机访问 自动验证并在训练期间显示验证数据负对数似然 支持自定义词二叉树构造。 使用由任何其他无监督模型构建的自己的二叉树来提高模型性能 更新 0.1 利用局部性优化GPU内存IO效率 使用 cuda...
CUDA全局内存的合并访问(个人理解) 每个warp去访问全局内存,会有400-600个时钟周期的内存延迟,这个代价很昂贵,所以为了减少访问全局内存的指令次数,我们将满足字节大小和对齐要求的warp合并起来访问全局内存,...
神经网络CUDA示例 神经网络工具包(PyTorch,TensorFlow等)的几个简单示例,这些示例调用自定义CUDA运算符。 我们提供了几种编译CUDA内核及其cpp包装程序的方法,包括jit,setuptools和cmake。 我们还提供了一些...
涵盖CUDA的指南,包括使您成为更好,更高效的CUDA开发人员的应用程序和工具。 注意:使用此方便的扩展程序,您可以轻松地将此markdown文件转换为PDF。 目录 CUDA工具包。 来源: CUDA学习资源 是NVIDIA开发的一种...
路径追踪 SS 这是一个用 CUDA C/C++ 编写的简单路径跟踪器它目前具有: 具有下一个事件估计的迭代路径跟踪光线的流合成重要性采样 Lambert、Phong 和 Cook-Torrance 着色器通过抖动抗锯齿区域和点光源反射、折射和...
JIT 编译流管理内核执行设备内存读/写不支持的功能内存池统一寻址事件和流事件流状态轮询流图捕获流式批处理内存操作外部存储器多设备助手(已经可能,但并不容易)图表纹理和表面OpenGL/VDPAU/EGL 互操作性例子有关...
GPU编程目前最流行,最可能成为国际标准的就是NVIDIA公司推出的CUDA编程模型。这个教程正是为众多GPU编程爱好者推出的,都是由国外的知名研究学者讲授,助你更深刻地理解GPU的体系结构,并对GPU程序进行优化
CUDA by example代码实例
cuda7.0:新功能实现主机多线程多流之间 实现kernel并行。主机的每个线程分配一个流 采用这个新特性可以实现多流之间的kernel并发执行
H.264编解码的CUDA实现,并行加速算法,内部实现了H.264视频编解码的CUDA实现,能够编译通过,已经过测试,并且还添加了注释信息。
论文研究-交通流分布式并行模拟的同步算法研究.pdf, 在对保守的、避免死锁的同步策略进行深入研究的基础上 ,提出了适用于交通流分布式系统预计费用估计和事件状态分析的...
第二部分是CUDA编程细节,对CUDA进行全方位的描述,包括内存,流和事件,执行模型(包括动态并行特性以及CUDA 5.0和SM 3.5的新特性),流处理器簇(包括SM 3.5的所有功能介绍),多GPU编程,纹理操作。这部分附带...
第1章从宏观上介绍流处理器演变历史。第2章详解GPU并行机制,深入理解串行与并行程序,以辩证地求解问题。第3章讲解CUDA设备及相关的硬件和体系结构,以实现优CUDA程序性能。第4章介绍CUDA开发环境搭建和可用调试...
用cuda实现的宽度优先搜索,但是由于cuda本身的原因,所读入的数据大于几万个时速度将比cpu快将近10倍。