CUDA By Examples 10 - Streams

最新推荐文章于 2023-03-01 12:01:26 发布
最新推荐文章于 2023-03-01 12:01:26 发布
阅读量623 收藏 1
点赞数
分类专栏： CUDA CUDA 编程 文章标签： cuda
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/full_speed_turbo/article/details/74332328
版权

一. Page-locked Host Memory

  1. 使用cudaHostAlloc()为Host分配内存属于页锁定内存, 这种内存不会被OS转移到disk上, 只能存在于物理内存中. 访问更安全. 但是只使用物理内存会使内存占用比例升高, 影响程序性能.
  2. 使用Page-locked Host Memory可以提高Host和GPU间的数据传输速度.
  3. 多用在cudaMemcpy的src或者dst上.
  4. Page-locked Host Memory要及时使用cudaFreeHost释放.

用不同内存分配方式测试数据传输速度:
(辅助代码见: http://blog.csdn.net/full_speed_turbo/article/details/71107132)

#include "../common/book.h"

#define SIZE    (100*1024*1024)

float cuda_malloc_test( int size, bool up)
{
    cudaEvent_t start, stop;
    int *a, *dev_a;
    float elapsedTime;

    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );

    HANDLE_ERROR( cudaMalloc( (void**)&dev_a, size*sizeof(int) ) );
    a = (int*)malloc(size * sizeof(int) );
    HANDLE_NULL(a);

    HANDLE_ERROR( cudaEventRecord( start, 0) );

    for (int i=0; i<100; i++)
    {
        if (up)
        {
            HANDLE_ERROR( cudaMemcpy( dev_a, a, size*sizeof(int), cudaMemcpyHostToDevice) );
        } 
        else
        {
            HANDLE_ERROR( cudaMemcpy( a, dev_a, size*sizeof(int), cudaMemcpyDeviceToHost ) );
        }
    }

    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime, start, stop ) );

    HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );
    HANDLE_ERROR( cudaFree(dev_a) );
    free(a);

    return elapsedTime;
}

float cuda_alloc_host_test(int size, bool up)
{
    cudaEvent_t start, stop;
    int *a, *dev_a;
    float elapsedTime;

    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );

    HANDLE_ERROR( cudaMalloc( (void**)&dev_a, size*sizeof(int) ) );
    HANDLE_ERROR( cudaHostAlloc( (void**)&a, size *sizeof(int), cudaHostAllocDefault ) );

    HANDLE_ERROR( cudaEventRecord( start, 0 ) );

    for (int i=0; i<100; i++)
    {
        if (up)
        {
            HANDLE_ERROR( cudaMemcpy(dev_a, a, size*sizeof(int), cudaMemcpyHostToDevice ) );
        } 
        else
        {
            HANDLE_ERROR( cudaMemcpy(a, dev_a, size*sizeof(int), cudaMemcpyDeviceToHost ) );
        }
    }

    HANDLE_ERROR( cudaEventRecord( stop, 0 ) );
    HANDLE_ERROR( cudaEventSynchronize( stop ) );
    HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime, start, stop ) );

    HANDLE_ERROR( cudaFreeHost( a ) );
    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );

    return elapsedTime;
}

int main(void)
{
    float elapsedTime;
    float MB = (float)100*SIZE*sizeof(int)/1024/1024;
    elapsedTime = cuda_malloc_test( SIZE, true );
    printf( "Time using malloc: %3.1f ms\n", elapsedTime );
    printf( "\t MB/s during copy up: %3.1f\n", MB/(elapsedTime/1000) );

    elapsedTime = cuda_malloc_test( SIZE, false );
    printf( "Time using malloc: %3.1f ms\n", elapsedTime );
    printf( "\t MB/s during copy down: %3.1f\n", MB/(elapsedTime/1000) );

    elapsedTime = cuda_alloc_host_test( SIZE, true );
    printf( "Time using HostAlloc: %3.1f ms\n", elapsedTime );
    printf( "\t MB/s during copy up: %3.1f\n", MB/(elapsedTime/1000) );

    elapsedTime = cuda_alloc_host_test( SIZE, false );
    printf( "Time using HostAlloc: %3.1f ms\n", elapsedTime );
    printf( "\t MB/s during copy down: %3.1f\n", MB/(elapsedTime/1000) );

    return 0;

}

输出:
这里写图片描述

2. streams

  1. stream可以看做GPU的一个task队列.
  2. 使用的Host内存需要用cudaHostAlloc分配.
  3. stream操作需要cudaStreamSynchronize同步.
  4. 多个stream配合时, 安排任务时要使用广度优先准则. 不要一次性把一个stream安排满了再去安排另外的stream. 否则, 另外的stream会被阻塞.

使用深度优先准则安排stream工作队列:

#include "../common/book.h"

#define N (1024*1024)
#define FULL_DATA_SIZE (N*20)

__global__ void kernel( 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(void)
{
    cudaDeviceProp prop;
    int whichDevice;
    HANDLE_ERROR( cudaGetDevice( &whichDevice ) );
    HANDLE_ERROR( cudaGetDeviceProperties( &prop, whichDevice ) );

    if(!prop.deviceOverlap)
    {
        printf("Device will not handle overlaps, so no speed up from streams\n");
        return 0;
    }

    cudaEvent_t start, stop;
    float elapsedTime;

    //开启计时器
    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );

    //初始化streams
    cudaStream_t stream0, 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();
    }

    //分段计算
    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
    for (int i=0; i<FULL_DATA_SIZE; i+=N*2)
    {
        //先安排stream0.
        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 ) );
        //安排stream1.
        HANDLE_ERROR( cudaMemcpyAsync( dev_a1, host_a+i+N,
                                        N * sizeof(int),
                                        cudaMemcpyHostToDevice,
                                        stream1 ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_b1, host_b+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 ) );
    }
    //同步streams.
    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 );

    //释放streams
    HANDLE_ERROR( cudaStreamDestroy( stream0 ) );
    HANDLE_ERROR( cudaStreamDestroy( stream1 ) );
    //释放Host内存
    HANDLE_ERROR( cudaFreeHost( host_a ) );
    HANDLE_ERROR( cudaFreeHost( host_b ) );
    HANDLE_ERROR( cudaFreeHost( host_c ) );
    //释放GPU内存
    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 ) );

    return 0;

}

这里写图片描述
使用广度优先原则安排stream任务队列:

    //分段计算
    HANDLE_ERROR( cudaEventRecord( start, 0 ) );
    for (int i=0; i<FULL_DATA_SIZE; i+=N*2)
    {
        //安排stream0.
        HANDLE_ERROR( cudaMemcpyAsync( dev_a0, host_a+i,
                                        N * sizeof(int),
                                        cudaMemcpyHostToDevice,
                                        stream0 ) );
        //安排stream1.
        HANDLE_ERROR( cudaMemcpyAsync( dev_a1, host_a+i+N,
                                        N * sizeof(int),
                                        cudaMemcpyHostToDevice,
                                        stream1 ) );
        //安排stream0.
        HANDLE_ERROR( cudaMemcpyAsync( dev_b0, host_b+i,
                                        N *sizeof(int),
                                        cudaMemcpyHostToDevice,
                                        stream0 ) );
        //安排stream1.
        HANDLE_ERROR( cudaMemcpyAsync( dev_b1, host_b+i+N,
                                        N * sizeof(int),
                                        cudaMemcpyHostToDevice,
                                        stream1 ) );
        //安排stream0.
        kernel<<<N/256, 256, 0, stream0>>>( dev_a0, dev_b0, dev_c0 );
        //安排stream1.
        kernel<<<N/256, 256, 0, stream1>>>( dev_a1, dev_b1, dev_c1 );
        //安排stream0.
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c0,
                                        N * sizeof(int),
                                        cudaMemcpyDeviceToHost,
                                        stream0 ) );
        //安排stream1.
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i+N, dev_c1,
                                        N*sizeof(int),
                                        cudaMemcpyDeviceToHost,
                                        stream1 ) );
    }

其他代码不变.
这里写图片描述

其实, 前期分配内存的时间比for循环里用的时间长的多.
所以, 对这个程序来说, 瓶颈在于Host和device之间的数据传输太慢了.

Full_Speed_Turbo
关注
  • 0
    点赞
  • 踩
  • 1
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
CUDA之Stream介绍
silencez_w的博客
12-10 6073
一、Stream的概念 用到CUDA的程序一般需要处理海量的数据，内存带宽经常会成为主要的瓶颈。在Stream的帮助下，CUDA程序可以有效地将内存读取和数值运算并行，从而提升数据的吞吐量。 Cuda stream是指一堆异步的cuda操作，他们按照host代码调用的顺序执行在device上。 典型的cuda编程模式我们已经熟知了： · 将输入数据从host转移到device · 在device上执行kernel · 将结果从device上转移回host Cuda Streams 所有的cuda操作（包括k
CUDA By Example(八)——流
hhhcbw的博客，欢迎各位来访
02-10 2133
CUDA 流操作，使用Stream加速程序执行
Cuda Streams的概述（一）-- Cuda介绍
大脸猫脸大主持人猫脸大的博客
05-14 1113
最近在做有关Cuda的一个项目， 碰到匪夷所思的问题，在异步的时候发现并没有达到预期的效果，程序没有异步起来，然后在网上找了一个Nvida的有关Cuda Streams的一个ppt，然后照着里面的提示，使程序达到了异步的效果。 首先，先回顾一下简单的简单的处理流程。 将数据从CPU内存中拷贝到GPU内存中。 启动一个GPU的kernel。 将结果从GPU内存中拷贝到CPU的内存中。 重复上面的操作。 在Pipelining中的并发相比于串型可以提高性能。 我们可以发现在并发的情况下，程序的运行速度有显
CUDA编程：stream和Concurrency初探
a135013563的博客
03-01 466
cuda
CUDA从入门到精通（七）：流并行
卜居
07-24 1万+
前面我们没有讲程序的结构，我想有些童鞋可能迫不及待想知道CUDA程序到底是怎么一个执行过程。好的，这一节在介绍流之前，先把CUDA程序结构简要说一下。 CUDA程序文件后缀为.cu，有些编译器可能不认识这个后缀的文件，我们可以在VS2008的Tools->Options->Text Editor->File Extension里添加cu后缀到VC++中，如下图：   一个.cu文件内既包
kafka-streams-examples:kafka-streams-example收集了我使用Kafka stream进行过的小型项目。 它的灵感来自kafka-src streams-examples
05-01
这些项目的灵感来自kafka-src streams-examples。 我计划也包括scala示例。 主题包括： WordCountLambdaDemo（工作示例）通过启动真正的Kafka服务器等来运行集成测试的指令位于该类的javadoc中 在流中运行...
go-streams:Go的轻量级流处理库
05-06
go-streams提供了一个简单而简洁的DSL来构建数据管道。 在计算中，管线（也称为数据管线）是一组串联连接的数据处理元素，其中一个元素的输出是下一个元素的输入。 流水线的元素通常以并行或按时间分割的方式执行...
akka-streams-kafka-examples
03-28
k8s-template的示例应用公开API端点： curl http://localhost:8080/api/hello/john记录INFO和ERROR日志。 名称为adam时出现错误，对于其他所有请求的INFOs： curl http://localhost:8080/api/hello/adam公开指标，...
akka-streams-kafka-examples-源码.rar
10-10
在"akka-streams-kafka-examples-源码.rar"压缩包中，我们可以深入理解Akka Streams如何与Kafka进行交互，以及如何利用它们的优势来处理大数据流。下面我们将详细探讨其中的关键知识点。 1. **Akka Streams基础** ...
kafka-streams-examples-0.10.2.1.jar
02-25
java运行依赖jar包
CUDA by example （中文：GPU高性能编程CUDA实战）代码实例
08-28
CUDA by example代码实例
cuda 函数前缀：device/global/host 相关问题
热门推荐
xiaoxiang_AQ 的博客
10-23 2万+
在深度学习caffe框架等多处，用到了CUDA函数，使代码加速，其文件类型为XXX.cu。在运行这些CUDA函数时，会遇到一下C++里没有的问题，比如函数前缀问题，如果在修改编写.cu文件时不注意，会出现问题，因此本文在这里讨论一下。主要解决如下问题：error : calling a __host__ function from a __global__ function is not al
CUDA + OpenCV实现CUDA By Example中波纹示例
小小学渣的博客
04-13 1060
参考了网上其他人的CUDA + OpenCV实现Julia Set的例子（ https://blog.csdn.net/Ricardo232525/article/details/74178963），用到了波纹示例上，代码如下：#include "device_launch_parameters.h"#include "cuda_runtime.h"  #include &lt;highgui.h...
NVIDIA CUDA 学习 (1) CUDA By Example
weixin_47658743的博客
09-05 375
CUDA By ExampleHost和DeviceCUDA C编程分配空间函数定义和执行使用规范 Host和Device host就是你的CPU，Device就是你的GPU。 CUDA C编程 #include <iostream> #include "book.h" __global__ void add( int a, int b, int *c ) { *c = a + b; } int main( void ) { int c; int *dev_c; HANDLE_ERROR
Cuda Streams的概述（六）-- 常见问题
大脸猫脸大主持人猫脸大的博客
05-19 761
常见问题 问题1:使用默认流 问题2：内存传输问题 问题3:隐式同步 问题4:主机端的限制 问题5:启动overhead限制 问题6:过度同步 问题7：PROFILER OVERHEAD case 1-A for(int i=0;i<repeat;i++) { kernel<<<1,1,0,stream1>>>(); kernel<<<1,1>>>(); } 存在问题：有一个kernel在默认流里。 解决方案：每个ker
CUDA By Examples 4 - 使用thread
Full_Speed_Turbo
06-26 578
使用block和thread
CUDA
u012129372的专栏
05-27 1409
CUDA是什么 CUDA，Compute Unified Device Architecture的简称，是由NVIDIA公司创立的基于他们公司生产的图形处理器GPUs（Graphics Processing Units,可以通俗的理解为显卡）的一个并行计算平台和编程模型。 通过CUDA，GPUs可以很方便地被用来进行通用计算（有点像在CPU中进行的数值计算等等）。在没有CUDA之前，GPUs一
CUDA By Example CUDA实战学习
answer_gusz的博客
11-13 515
刚开始学习GPU编程，看大家推荐CUDA By Example,刚开始看会觉得有些困惑，毕竟没啥基础知识，所以先找了本《高性能CUDA应用设计与开发》，看了其中几章对于GPU和CUDA编程的基础知识讲解，有了大概了解以后开始学习CUDA实战。毕竟高性能CUDA这本书没什么具体编程的东西，想提高还得上手练习。 在用CUDA By Example 这本书的时候，发现按照书本的例子，总是提示错误，主要
CUDA学习笔记 02 函数三种前缀device、global、host
JackZhangNJU的专栏
09-18 3309
CUDA三种前缀分别用于在定义函数时限定该函数的调用和执行方式，如下： （1）_host_ int foo(int a){}与C或者C++中的foo(int a){}相同，是由CPU调用，由CPU执行的函数 。 和我们平常写的函数是一致的。 __host__ int foo(int a) { return 1； } （2）_global_ int foo(int a){}表示一
cesium-examples-master
最新发布
07-25
### 回答1： cesium-examples-master 是一个 Cesium 的示例项目。Cesium 是一个开源的3D地球可视化引擎，能够在Web上以浏览器为平台展示地球相关的数据和图形。cesium-examples-master 包含了一系列基于 Cesium 引擎的示例代码和样例数据，供开发人员学习和参考。 这个项目提供了丰富的示例，涵盖了各种场景和功能，如地形渲染、卫星图像展示、空中飞行效果、地球热力图、数据可视化等。每个示例都提供了完整的源代码和相关资源，开发人员可以直接运行和修改，快速了解 Cesium 的使用方式和功能特性。 cesium-examples-master 的目的是帮助开发人员加快上手 Cesium，提供具体的示例代码和实现思路，同时也是一个社区贡献的项目，任何人都可以向其中添加自己的示例代码。这对于想要共享自己的 Cesium 开发经验，或者想要通过Cesium实现自己的创意项目的开发者们来说都是很有帮助的。 总之，cesium-examples-master 是一个集合了Cesium引擎的示例代码和样例数据的项目，通过这个项目，开发人员可以学习和参考Cesium的使用方式和功能特性，同时也可以贡献自己的示例代码，为Cesium社区贡献自己的力量。 ### 回答2： cesium-examples-master是一个开源的Cesium.js示例库。Cesium.js是一个基于WebGL的开源JavaScript库，用于创建3D地球和地理信息可视化应用程序。 cesium-examples-master库中包含了大量的示例代码，用于演示如何使用Cesium.js库进行地球和地理数据可视化。这些示例涵盖了各种应用场景，包括地球浏览、地理数据可视化、飞行模拟、地球时间轴等等。 这个示例库非常有用，特别是对于那些想要利用Cesium.js构建自己的3D地球和地理信息应用程序的开发人员来说。通过学习和理解这些示例代码，开发人员可以快速上手并加快应用程序的开发速度。 此外，cesium-examples-master还可以作为一个学习资源，供初学者学习Cesium.js库的使用。通过运行和修改这些示例代码，初学者可以逐步掌握Cesium.js的各种功能和技术知识。 总之，cesium-examples-master是一个非常有用的示例库，可以帮助开发人员和初学者更好地了解和应用Cesium.js库。无论是开发3D地球和地理信息应用程序，还是学习Cesium.js库的使用，这个示例库都是一个很好的资源。如果你对Cesium.js感兴趣，不妨去查看cesium-examples-master库并尝试运行其中的示例代码。 ### 回答3： cesium-examples-master是一个Cesium的示例代码库。Cesium是一个开源的地球可视化库，用于在Web浏览器中创建交互式三维地球和地球数据的应用程序。cesium-examples-master提供了许多不同类型的示例，展示了使用Cesium创建各种地球可视化应用的能力。这些示例包括地球模型的加载、地形数据的展示、地图投影的转换、地球上的点、线和面的创建等等。通过这些示例，开发者可以学习如何使用Cesium的API来构建自己的地球可视化项目，并根据自己的需求进行修改和扩展。cesium-examples-master的代码注释详细，对于刚开始学习Cesium的开发者来说是一个很好的参考工具。在cesium-examples-master中，开发者可以找到各种应用场景的示例代码，例如飞行模拟、地球上的图层切换、轨迹的绘制和动态效果等等。总之，cesium-examples-master对于想要学习和探索Cesium地球可视化库的开发者来说是一个非常有用的资源。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值