CUDA 学习——常量内存与事件（CUDA实战 第六章）

最新推荐文章于 2024-04-20 22:27:39 发布
最新推荐文章于 2024-04-20 22:27:39 发布
阅读量433 收藏 3
点赞数 1
分类专栏： CUDA实战 文章标签： 学习 c++ c语言
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/ygf666/article/details/124844567
版权

1.光线跟踪

#include <stdio.h>
#include <GL/glut.h> 
struct CPUBitmap {
    unsigned char    *pixels;
    int     x, y;
    void    *dataBlock;
    void (*bitmapExit)(void*);

    CPUBitmap( int width, int height, void *d = NULL ) {
        pixels = new unsigned char[width * height * 4];
        x = width;
        y = height;
        dataBlock = d;
    }

    ~CPUBitmap() {
        delete [] pixels;
    }

    unsigned char* get_ptr( void ) const   { return pixels; }
    long image_size( void ) const { return x * y * 4; }

    void display_and_exit( void(*e)(void*) = NULL ) {
        CPUBitmap**   bitmap = get_bitmap_ptr();
        *bitmap = this;
        bitmapExit = e;
        // a bug in the Windows GLUT implementation prevents us from
        // passing zero arguments to glutInit()
        int c=1;
        char* dummy = "";
        glutInit( &c, &dummy );
        glutInitDisplayMode( GLUT_SINGLE | GLUT_RGBA );
        glutInitWindowSize( x, y );
        glutCreateWindow( "bitmap" );
        glutKeyboardFunc(Key);
        glutDisplayFunc(Draw);
        glutMainLoop();
    }

     // static method used for glut callbacks
    static CPUBitmap** get_bitmap_ptr( void ) {
        static CPUBitmap   *gBitmap;
        return &gBitmap;
    }

   // static method used for glut callbacks
    static void Key(unsigned char key, int x, int y) {
        switch (key) {
            case 27:
                CPUBitmap*   bitmap = *(get_bitmap_ptr());
                if (bitmap->dataBlock != NULL && bitmap->bitmapExit != NULL)
                    bitmap->bitmapExit( bitmap->dataBlock );
                exit(0);
        }
    }

    // static method used for glut callbacks
    static void Draw( void ) {
        CPUBitmap*   bitmap = *(get_bitmap_ptr());
        glClearColor( 0.0, 0.0, 0.0, 1.0 );
        glClear( GL_COLOR_BUFFER_BIT );
        glDrawPixels( bitmap->x, bitmap->y, GL_RGBA, GL_UNSIGNED_BYTE, bitmap->pixels );
        glFlush();
    }
};

#define DIM 1024

#define rnd( x ) (x * rand() / RAND_MAX)
#define INF     2e10f

struct Sphere {
    float   r,b,g;
    float   radius;
    float   x,y,z;
    __device__ float hit( float ox, float oy, float *n ) {
        float dx = ox - x;
        float dy = oy - y;
        if (dx*dx + dy*dy < radius*radius) {
            float dz = sqrtf( radius*radius - dx*dx - dy*dy );
            *n = dz / sqrtf( radius * radius );
            return dz + z;
        }
        return -INF;
    }
};
#define SPHERES 20



__global__ void kernel( Sphere *s, unsigned char *ptr ) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;
    float   ox = (x - DIM/2);
    float   oy = (y - DIM/2);

    float   r=0, g=0, b=0;
    float   maxz = -INF;
    for(int i=0; i<SPHERES; i++) {
        float   n;
        float   t = s[i].hit( ox, oy, &n );
        if (t > maxz) {
            float fscale = n;
            r = s[i].r * fscale;
            g = s[i].g * fscale;
            b = s[i].b * fscale;
            maxz = t;
        }
    } 

    ptr[offset*4 + 0] = (int)(r * 255);
    ptr[offset*4 + 1] = (int)(g * 255);
    ptr[offset*4 + 2] = (int)(b * 255);
    ptr[offset*4 + 3] = 255;
}

// globals needed by the update routine
struct DataBlock {
    unsigned char   *dev_bitmap;
    Sphere          *s;
};

int main( void ) {
    
    DataBlock   data;
    cudaEvent_t     start, stop;
    cudaEventCreate( &start );
    cudaEventCreate( &stop );
    cudaEventRecord( start, 0 );

    CPUBitmap bitmap( DIM, DIM, &data );
    unsigned char   *dev_bitmap;

    // allocate memory on the GPU for the output bitmap
    cudaMalloc( (void**)&dev_bitmap,
                              bitmap.image_size() );
    Sphere *s;
    cudaMalloc( (void**)&s,
                      sizeof(Sphere) * SPHERES );

    // allocate temp memory, initialize it, copy to constant
    // memory on the GPU, then free our temp memory
    Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere) * SPHERES );
    for (int i=0; i<SPHERES; i++) {
        temp_s[i].r = rnd( 1.0f );
        temp_s[i].g = rnd( 1.0f );
        temp_s[i].b = rnd( 1.0f );
        temp_s[i].x = rnd( 1000.0f ) - 500;
        temp_s[i].y = rnd( 1000.0f ) - 500;
        temp_s[i].z = rnd( 1000.0f ) - 500;
        temp_s[i].radius = rnd( 100.0f ) + 20;
    }
    
    cudaMemcpy( s, temp_s,
        sizeof(Sphere) * SPHERES,
        cudaMemcpyHostToDevice );
    free( temp_s );

    // generate a bitmap from our sphere data
    dim3    grids(DIM/16,DIM/16);
    dim3    threads(16,16);
    kernel<<<grids,threads>>>( s, dev_bitmap );

    // copy our bitmap back from the GPU for display
    cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
                              bitmap.image_size(),
                              cudaMemcpyDeviceToHost );

    // get stop time, and display the timing results
    cudaEventRecord( stop, 0 );
    cudaEventSynchronize( stop );
    float   elapsedTime;
    cudaEventElapsedTime( &elapsedTime,
                                        start, stop ) ;
    printf( "Time to generate:  %3.1f ms\n", elapsedTime );

    cudaEventDestroy( start ) ;
    cudaEventDestroy( stop ) ;

    cudaFree( dev_bitmap ) ;
    cudaFree( s );

    // display
    bitmap.display_and_exit();
}

编译，输出：
用时4+ms
在这里插入图片描述

2.常量内存实现光线跟踪

#include <stdio.h>
#include <GL/glut.h> 
struct CPUBitmap {
    unsigned char    *pixels;
    int     x, y;
    void    *dataBlock;
    void (*bitmapExit)(void*);

    CPUBitmap( int width, int height, void *d = NULL ) {
        pixels = new unsigned char[width * height * 4];
        x = width;
        y = height;
        dataBlock = d;
    }

    ~CPUBitmap() {
        delete [] pixels;
    }

    unsigned char* get_ptr( void ) const   { return pixels; }
    long image_size( void ) const { return x * y * 4; }

    void display_and_exit( void(*e)(void*) = NULL ) {
        CPUBitmap**   bitmap = get_bitmap_ptr();
        *bitmap = this;
        bitmapExit = e;
        // a bug in the Windows GLUT implementation prevents us from
        // passing zero arguments to glutInit()
        int c=1;
        char* dummy = "";
        glutInit( &c, &dummy );
        glutInitDisplayMode( GLUT_SINGLE | GLUT_RGBA );
        glutInitWindowSize( x, y );
        glutCreateWindow( "bitmap" );
        glutKeyboardFunc(Key);
        glutDisplayFunc(Draw);
        glutMainLoop();
    }

     // static method used for glut callbacks
    static CPUBitmap** get_bitmap_ptr( void ) {
        static CPUBitmap   *gBitmap;
        return &gBitmap;
    }

   // static method used for glut callbacks
    static void Key(unsigned char key, int x, int y) {
        switch (key) {
            case 27:
                CPUBitmap*   bitmap = *(get_bitmap_ptr());
                if (bitmap->dataBlock != NULL && bitmap->bitmapExit != NULL)
                    bitmap->bitmapExit( bitmap->dataBlock );
                exit(0);
        }
    }

    // static method used for glut callbacks
    static void Draw( void ) {
        CPUBitmap*   bitmap = *(get_bitmap_ptr());
        glClearColor( 0.0, 0.0, 0.0, 1.0 );
        glClear( GL_COLOR_BUFFER_BIT );
        glDrawPixels( bitmap->x, bitmap->y, GL_RGBA, GL_UNSIGNED_BYTE, bitmap->pixels );
        glFlush();
    }
};

#define DIM 1024

#define rnd( x ) (x * rand() / RAND_MAX)
#define INF     2e10f

struct Sphere {
    float   r,b,g;
    float   radius;
    float   x,y,z;
    __device__ float hit( float ox, float oy, float *n ) {
        float dx = ox - x;
        float dy = oy - y;
        if (dx*dx + dy*dy < radius*radius) {
            float dz = sqrtf( radius*radius - dx*dx - dy*dy );
            *n = dz / sqrtf( radius * radius );
            return dz + z;
        }
        return -INF;
    }
};
#define SPHERES 20
//下面有修改
__constant__ Sphere s[SPHERES];
//下面有修改
__global__ void kernel( unsigned char *ptr ) {
    // map from threadIdx/BlockIdx to pixel position
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;
    float   ox = (x - DIM/2);
    float   oy = (y - DIM/2);

    float   r=0, g=0, b=0;
    float   maxz = -INF;
    for(int i=0; i<SPHERES; i++) {
        float   n;
        float   t = s[i].hit( ox, oy, &n );//计算来自（ox，oy）的光线是否与这个球面相交
        if (t > maxz) {
            float fscale = n;
            r = s[i].r * fscale;
            g = s[i].g * fscale;
            b = s[i].b * fscale;
            maxz = t;
        }
    } 

    ptr[offset*4 + 0] = (int)(r * 255);
    ptr[offset*4 + 1] = (int)(g * 255);
    ptr[offset*4 + 2] = (int)(b * 255);
    ptr[offset*4 + 3] = 255;
}

// globals needed by the update routine
struct DataBlock {
    unsigned char   *dev_bitmap;
};

int main( void ) {
    DataBlock   data;
    // capture the start time
    cudaEvent_t     start, stop;
    cudaEventCreate( &start );
    cudaEventCreate( &stop );
    cudaEventRecord( start, 0 );

    CPUBitmap bitmap( DIM, DIM, &data );
    unsigned char   *dev_bitmap;

    // allocate memory on the GPU for the output bitmap
    cudaMalloc( (void**)&dev_bitmap,
                              bitmap.image_size() );

    // allocate temp memory, initialize it, copy to constant
    // memory on the GPU, then free our temp memory
    Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere) * SPHERES );
    for (int i=0; i<SPHERES; i++) {
        temp_s[i].r = rnd( 1.0f );
        temp_s[i].g = rnd( 1.0f );
        temp_s[i].b = rnd( 1.0f );
        temp_s[i].x = rnd( 1000.0f ) - 500;
        temp_s[i].y = rnd( 1000.0f ) - 500;
        temp_s[i].z = rnd( 1000.0f ) - 500;
        temp_s[i].radius = rnd( 100.0f ) + 20;
    }
    //下面有修改
    cudaMemcpyToSymbol( s, temp_s, sizeof(Sphere) * SPHERES);
    free( temp_s );

    // generate a bitmap from our sphere data
    dim3    grids(DIM/16,DIM/16);
    dim3    threads(16,16);
    kernel<<<grids,threads>>>( dev_bitmap );

    // copy our bitmap back from the GPU for display
    cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
                              bitmap.image_size(),
                              cudaMemcpyDeviceToHost );

    // get stop time, and display the timing results
    cudaEventRecord( stop, 0 );
    cudaEventSynchronize( stop );
    float   elapsedTime;
    cudaEventElapsedTime( &elapsedTime,
                                        start, stop ) ;
    printf( "Time to generate:  %3.1f ms\n", elapsedTime );

    cudaEventDestroy( start ) ;
    cudaEventDestroy( stop ) ;

    cudaFree( dev_bitmap ) ;

    // display
    bitmap.display_and_exit();
}

编译，输出：花费时间3-4ms左右
在这里插入图片描述

3.总结与分析

总的看，加速的效果一般，增加球的数量try try

修改为10000个球，非常量内存结果为：837.8ms
在这里插入图片描述
常量内存 发现10000个球放不下！
改为1000个，结果：76.9ms
在这里插入图片描述
再把非常量内存也改为1000,结果：99.7ms
在这里插入图片描述

风尘23187
关注
  • 1
    点赞
  • 踩
  • 3
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 0
    评论
专栏目录
CUDA进阶：内存模型
Shilong的博客
09-19 1140
CUDA内存模型
CUDA_常量内存
weixin_48151686的博客
11-09 719
常量内存简介 常量内存其实只是全局内存的一种虚拟地址形式，并没有特殊保留的常量内存块。常量内存有两个特性： 高速缓存,拥有缓存加速； 支持将单个值广播到线程束中的每个线程 常量内存的大小限定为64K，每个SM拥有8KB的常数存储器缓存，在编译期时声明一块常量内存，需要用到__constant__关键字，例如 __constant__ float my_array[1024] = { 0.0F, 1.0F, 1.3F, ...}; 不同于c/c++中的const常量，cuda中常量内存在声明后是可以修改
GPU高性能编程CUDA实战—book.h CPUBitmap.h等头文件
04-09
GPU高性能编程CUDA实战—book.h CPUBitmap.h等头文件
CUDA编程---共享内存和常量内存
最新发布
qq_40268672的博客
04-20 822
是较大的板载内存，具有相对较高的延迟。是较小的片上内存，具有相对较低的延迟，并且共享内存可以提供比全局内存高得多的带宽。可以把它当作一个。
NVIDIA CUDA2023春训营（八）CUDA 事件
Alex_McAvoy的博客
02-07 764
NVIDIA CUDA2023春训营
光线追踪算法
爱学习的小刘芒的博客
06-24 5002
光线追踪算法
cuda编程学习——CUDA内存介绍（七）
出门吃三碗饭的博客
05-31 1274
现代计算机中的内存往往存在一种组织结构（hierarchy）。在这种结构中，含有多种类 型的内存，每种内存分别具有不同的容量和延迟（latency，可以理解为处理器等待内存数据的时间）。一般来说，延迟低（速度高）的内存容量小，延迟高（速度低）的内存容量大。
CUDA-F-5-5-常量内存
Tony的博客
06-26 368
title: 【CUDA 基础】5.5 常量内存 categories: - CUDA - Freshman tags: - CUDA常量内存 - CUDA只读缓存 toc: true date: 2018-06-06 06:20:07 Abstract: 本文介绍另外两种内存——常量内存，只读缓存 Keywords: CUDA常量内存，CU...
[cuda]常量内存
手上有剑，心中有爱
12-14 230
    我们在cuda编程中定义变量的时候，可以用__constant__来修饰定义，代表着定义了一个常量内存。向常量内存中拷贝数据的时候不再使用cudaMemcpy，而改为使用cudaMemcpyToSymbol。对于常量内存，不需要在使用完成后用cudafree释放空间。     常量内存带来的性能提升     __constant__把对变量的访问限制为只读，跟从全局内存读取数据相比...
CUDA常量内存
csdn_hao的专栏
09-06 297
cuda中常量数组的使用：
CUDA中的常量内存__constant__
热门推荐
牧野的博客
02-06 1万+
GPU包含数百个数学计算单元，具有强大的处理运算能力，可以强大到计算速率高于输入数据的速率，即充分利用带宽，满负荷向GPU传输数据还不够它计算的。CUDA C除全局内存和共享内存外，还支持常量内存，常量内存用于保存在核函数执行期间不会发生变化的数据，使用常量内存在一些情况下，能有效减少内存带宽，降低GPU运算单元的空闲等待。 使用常量内存提升性能 使用常量内存可以提升运算性能的
cuda立体匹配——中值滤波
06-20
匹配代价计算的目的是衡量待匹配像素与候选像素之间的相关性。两个像素无论是否为同名点，都可以通过匹配代价函数计算匹配代价，代价越小则说明相关性越大，是同名点的概率也越大。匹配代价计算的方法有很多，本课设...
CUDA——性能优化（一）
01-20
CUDA全局内存的合并访问（个人理解） 每个warp去访问全局内存，会有400-600个时钟周期的内存延迟，这个代价很昂贵，所以为了减少访问全局内存的指令次数，我们将满足字节大小和对齐要求的warp合并起来访问全局内存，...
cuda.rar_cuda_cuda学习_cuda编程
09-24
CUDA学习入门教程详细讲解CUDA编程方法，适合初学者学习。
cuda-handbook.rar_cuda_cuda学习
09-20
CUDA编程学习入门笔记，很不错。可以知道cuda是干什么用的
cuda-example.rar_cuda examples_cuda 例子_cuda学习
09-19
cuda examples 对于学习GPU应用的人有用，例子讲的很详细具体
CUDA内存管理一文理清
HinGwenWoong的博客
02-13 1507
CUDA内存管理
CUDA By Example笔记--常量内存与事件
牵一只蜗牛去散步
04-12 167
当处理常量内存时，NVIDIA硬件将单次内存读取操作广播到半线程束中（16个线程）；当半线程束的每个线程都从常量内存相同地址读取数据时，GPU只会产生一次读取请求并将数据广播到每个线程中；因此，当从常量内存中读取大量数据时，产生的内存流量仅为使用全局变量的1/16；对常量内存的单次读操作可以广播到其他临近的线程，节约临近线程的读取操作；常量内存的数据将被缓存，对相同地址的连续读操作不会产生额外的内存通信量；CUDA事件的本质是一个 GPU 时间戳，其在用户指定的时间点上被记录；
基于CUDA的光线跟踪算法
BugRunner的专栏
09-16 1万+
光线跟踪是目前常用的主流绘制技术之一，由于它能方便地模拟生成复杂的光照效果，生成高质量的图像，在很多领域得到了广泛的应用，如真实感绘制，虚拟现实，可视化，计算机动画等。但光线跟踪算法的计算开销较大，妨碍了其应用效率。光线跟踪渲染的主要运算操作耗费在光线与场景求交的过程中，为此，为了加快这一计算过程，人们研究了大量的技术来加快求交操作。其中，通过建立一定的空间组织结构来加速求交计算机取得了很大的进展
关于cuda常量内存与全局内存
04-23
cuda常量内存和全局内存都是GPU中存储数据的区域，但它们之间有一些区别。常量内存的数据是被编译时确定的常量，无法被修改，且常量内存的读取速度相对全局内存更快。而全局内存则是数据可被修改的变量存储区域，它...

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

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

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

风尘23187

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

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

余额充值