第6章常量内存与事件

Constant Memory（常量内存），是一种内存区域事件：用来测量CUDA应用程序的性能，用来定量的分析对应用程序的某个修改是否带来性能提升。

6.1 本章目标

了解如何在CUDA C中使用常量内存。
了解常量内存的性能特性。
学习如何使用CUDA事件来测量应用程序的性能

6.2 常量内存

用于保存在核函数执行期间不会发生变化的数据。NVIDIA硬件提供了64KB的常量内存，并且对常量内存采取了不同标准全局内存的处理方式。在某些情况中用常量内存来替换全局内存能有效地减少内存带宽。

6.2.1 光纤跟踪简介

光纤跟踪（Ray Tracing）示例介绍如何使用常量内存。

背景：光线跟踪是一种真实地显示物体的方法，该方法由Appe在1968年提出。光线跟踪方法沿着到达视点的光线的反方向跟踪，经过屏幕上每一个象素，找出与视线相交的物体表面点P0，并继续跟踪，找出影响P0点光强的所有光源，从而算出P0点上精确的光线强度，在材质编辑中经常用来表现镜面效果。光线跟踪或称光迹追踪是计算机图形学的核心算法之一。在算法中，光线从光源被抛射出来，当他们经过物体表面的时候，对他们应用种种符合物理光学定律的变换。最终，光线进入虚拟的摄像机底片中，图片被生成出来。

简单的说，光纤跟踪是从三维对象场景中生成二维图像的一种方式。原理，在场景中选择一个位置放上一台假想的相机。这台数字相机包含一个光传感器来生成图像，因此我们需要判断哪些光接触到这个传感器。图像中的每个像素与命中传感器的光线有着相同的颜色和强度。

光纤跟踪实现了哪些功能：它将从每个像素发射一道光线，并且跟踪这些光线会命中哪些球面。此外，它还将跟踪每道命中光线的深度。当一道光线穿过多个球面时，只有最接近相机的球面才会被看到。我们的“光纤跟踪器”会把相机看不到的球面隐藏起来。

6.2.2 通过普通方式实现光线跟踪

通过一个数据结果对球面建模，在数据结构中包含了球面的中心坐标(x, y, z)，半径radius，以及颜色值(r, g, b)。

#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;
    }
};

hit( float ox, float oy, float *n )，对于来自(ox, oy) 处像素的光线，这个方法将计算光线是否与这个球面相交，那么这个方法将计算从相机到光线命中球面处的距离。当一道光线穿过多个球面时，只有最接近相机的球面才会被看到。

main() 代码结构

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

#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;
    // capture the start time
    cudaEvent_t     start, stop;
    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );
    HANDLE_ERROR( cudaEventRecord( start, 0 ) );

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


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

    // allocate temp memory, initialize it, copy to
    // 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;
    }
    HANDLE_ERROR( 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
    HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
                              bitmap.image_size(),
                              cudaMemcpyDeviceToHost ) );

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

    HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );

    HANDLE_ERROR( cudaFree( dev_bitmap ) );
    HANDLE_ERROR( cudaFree( s ) );

    // display
    bitmap.display_and_exit();
}

6.2.3 通过常量内存来实现光线跟踪

这个示例中只有一个输入数据，即球面数组，因此可以将这个数据保存到常量内存中。

声明修饰符：__constant__ 例如：__constant__ Sphere s[SPHERES];

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

#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 );
        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;
    HANDLE_ERROR( cudaEventCreate( &start ) );
    HANDLE_ERROR( cudaEventCreate( &stop ) );
    HANDLE_ERROR( cudaEventRecord( start, 0 ) );

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

    // allocate memory on the GPU for the output bitmap
    HANDLE_ERROR( 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;
    }
    HANDLE_ERROR( 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
    HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
                              bitmap.image_size(),
                              cudaMemcpyDeviceToHost ) );

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

    HANDLE_ERROR( cudaEventDestroy( start ) );
    HANDLE_ERROR( cudaEventDestroy( stop ) );

    HANDLE_ERROR( cudaFree( dev_bitmap ) );

    // display
    bitmap.display_and_exit();
}

修改：

HANDLE_ERROR( cudaMemcpyToSymbol( s, temp_s, 
                                sizeof(Sphere) * SPHERES) );

cudaMemcpyToSymbol() 会复制到常量内存，而cudaMemcpy() 会复制到全局内存

6.2.4 常量内存带来的性能提升

__constant__ 将把变量的访问限制为只读，且节约内存带宽

对常量内存的单次读操作可以广播到其他的“临近(Nearby)”线程，这将节约15次读取操作。
常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。

Warp: 线程束，可以看成是一组线程通过交织而形成的一个整体。在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。在程序中的每一行，线程束中的每个线程都将在不同的数据上执行相同的指令。

当处理常量内存时，NVIDIA硬件将把单次内存读取操作广播到每个半线程束（Half-Warp）。在半线程束中包含了16个线程，即线程束中线程数量的一半。如果在半线程束中的每个线程都从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量的数据，那么这种方式产生内存流量只是使用全局内存时的1/16（大约6%）。

但在读取常量内存时，所节约的并不仅限于减少了94%的带宽。由于这块内存的内容是不会发生变化的，因此硬件将主动将这个常量的数据缓存在GPU上。在第一次从常量内存的某个地址上读取后，当其他半线程束请求同一个地址时，那么命中缓存，这同样减少了额外的内存流量。

在我们的光线跟踪器中，每个线程都读取球面的相应数据而从计算它与光线的相交情况。在把应用程序改为将球面数据保存在常量内存后，硬件只需要请求这个数据一次。在缓存数据后，其他每个线程将不会产生内存流量，原因：

线程将在半线程束的广播中收到这个数据。
从常量内存缓存中收到数据。

然而，在使用常量内存时，也可能对性能产生负面影响。半线程束广播功能实际上是一把双刃剑。虽然当所有16个线程都读取相同地址时，这个功能可以极大的提升性能，但当所有16个线程分别读取不同地址时，它实际上会降低性能。

只有当16个线程每次都需要相同的读取请求时，才值得将这个读取操作广播到16个线程。然而，如果半线程束中所有线程需要访问常量内存中不同的数据，那么这个16次不同的读取操作会被串行化，从而需要16的时间来发出请求。但如果从全局内存中读取，那么这些请求会同时发出。在这种情况中，从常量内存读取就慢于从全局内存中读取。

6.3 使用事件来测量性能

CUDA 中的事件本质上是一个GPU时间戳。

<pre name="code" class="cpp">// 创建事件： 
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// 记录开始事件：
cudaEventRecord(start, o);

// 在GPU上执行一些工作
// ...
// 记录结束事件：
cudaEventRecord(&stop, 0);
// 事件同步 
cudaEventSynchronize(stop);

测量光线跟踪器的性能 cudaEventElapsedTime( &elapsedTime,

start, stop ) );

是一个工具函数，用来计算两个事件之间经历的时间。第一个参数：某个浮点变量地址，在这个参数中将包含两次事件之间经历的时间，单位为毫秒。

<pre name="code" class="cpp">cudaEventDestroy( stop ); // 是一个销毁函数。

《GPU高性能编程CUDA实战》学习笔记（六）第6章常量内存与事件

第6章常量内存与事件

6.1 本章目标

6.2 常量内存

6.2.1 光纤跟踪简介

6.2.2 通过普通方式实现光线跟踪

6.2.3 通过常量内存来实现光线跟踪

6.2.4 常量内存带来的性能提升

6.3 使用事件来测量性能

继续阅读

（C# 编程指南）

java 开发中的骆驼拼写法（CamelCase）

20/20: Top 20 Programming Lessons I've Learned in 20 Years

20条编程经验

【观点】风雨20年：我所积累的20条编程经验【观点】风雨20年：我所积累的20条编程经验

【转载】风雨20年：我所积累的20条编程经验

如果你的梦想是成为一名程序员，那可能需要看看这篇文章！

ImportError: libcublas.so.9.0: cannot open shared object file: No such file or directory（完美解决）

一种解决思路： ImportError: libcublas.so.10.0: cannot open shared object file: No such file

写在工作13周年

编程高手：VC 6.0下利用消息实现内部进程通讯

VS2008 不在支持 Attributed ATL

C++之父Bjarne Stroustrup的演讲：软件理想和历史

linux网络编程----发送与接收文件

各种二分查找

JAVA高效编程指南

《GPU高性能编程CUDA实战》学习笔记（六）第6章 常量内存与事件

第6章 常量内存与事件

6.1 本章目标

6.2 常量内存

6.2.1 光纤跟踪简介

6.2.2 通过普通方式实现光线跟踪

6.2.3 通过常量内存来实现光线跟踪

6.2.4 常量内存带来的性能提升

6.3 使用事件来测量性能

继续阅读

《GPU高性能编程CUDA实战》学习笔记（六）第6章常量内存与事件

第6章常量内存与事件