天天看點
傳回

《GPU高性能程式設計CUDA實戰》學習筆記（六）第6章 常量記憶體與事件

第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 ); // 是一個銷毀函數。
cuda 顯示卡 高性能 gpu 程式設計
上一篇: cuda矩陣之心得矩陣程式設計，cuda程式設計入門
下一篇: 《GPU高性能程式設計CUDA實戰》學習筆記（七）第七章 紋理記憶體

繼續閱讀