GPU程式設計自學7 常量記憶體與事件

gpu通常包含大量的數學計算單元，因此效能瓶頸往往不在於晶元的數學計算吞吐量，而在於晶元的記憶體頻寬，即有時候輸入資料的速率甚至不能維持滿負荷的運算。於是我們需要一些手段來減少記憶體通訊量。目前的gpu均提供了64kb的常量記憶體，並且對常量記憶體採取了不同於全域性記憶體的處理方式。在某些場景下，使用常量記憶體來替換全域性記憶體可以有效地提高通訊效率。

7.1 常量記憶體

常量記憶體具有以下特點：

常量記憶體帶來效能提公升的原因主要有兩個：

對於原因1，涉及到執行緒束(warp)的概念。

在cuda架構中，執行緒束是指乙個包含32個執行緒的集合，這個執行緒集合被「編織在一起」並且以「步調一致(lockstep)」的形式執行。即執行緒束中的每個執行緒都將在不同資料上執行相同的指令。

當處理常量記憶體時，nvidia硬體將把單次記憶體讀取操作廣播到每個半執行緒束(half-warp)。在半執行緒束中包含16個執行緒，即執行緒束中線程數量的一半。如果在半執行緒束中的每個執行緒從常量記憶體的相同位址上讀取資料，那麼gpu只會產生一次讀取請求並在隨後將資料廣播到每個執行緒。如果從常量記憶體中讀取大量資料，那麼這種方式產生的記憶體流量只是使用全域性記憶體時的1/16。

對於原因2，涉及到快取的管理

由於常量記憶體的內容是不發生變化的，因此硬體將主動把這個常量資料快取在gpu上。在第一次從常量記憶體的某個位址上讀取後，當其他半執行緒束請求同乙個位址時，那麼將命中快取，這同樣減少了額外的記憶體流量。

另一方面, 常量記憶體的使用也可能會對效能產生負面的影響。半執行緒束廣播功能實際上是一把雙刃劍。雖然當所有16個執行緒都讀取相同位址時，這個功能可以極大提公升效能，但當所有16個執行緒分別讀取不同的位址時，它實際上會降低效能。因為這16次不同的讀取操作會被序列化，從而需要16倍的時間來發出請求。但如果從全域性記憶體中讀取，那麼這些請求會同時發出。

7.2 常量記憶體應用例項 —— 光線跟蹤

下面通過乙個光線跟蹤的例項來說明一下常量記憶體的使用效果。

下面的光線跟蹤不涉及光源以及光線反射，只是簡單的類似於「投影」的操作，如下圖所示。

投影平面前面會有大量存在重疊的球體(這裡我只畫了1個)，投影平面上每個畫素點會發射出一條射線(射線方向認為是z方向)，我們需要和所有球體判斷相交情況。如果和多個球體相交，則選擇最近的交點(即無法看到遮擋的球體)。根據交點到對應球心的距離(z方向距離)確定投影點的畫素值，距離越遠則越亮。

如果距離為無窮大，則表明沒有相交，則置為黑色背景。

由於每個畫素都會射出一條射線，然後和所有球體計算相交，因此需要經常訪問固定的球體引數。因此，為了提高訪問效率，我們將球體資訊定義到常量記憶體。

**如下(需要opencv)：

#include "cuda_runtime.h"  
#include "highgui.hpp"  
#include 
using
namespace cv;
#define inf 2e10f // 定義無窮遠距離(用於表示沒有球體相交) 
#define rnd(x) (x*rand()/rand_max) 
#define spheres 100   //球體數量
#define dim 1024     //影象大小
// 球體資訊結構體
struct sphere
return -inf;
}};// 宣告球體陣列
__constant__ sphere s[spheres];
// 光線跟蹤核函式
//__global__ void raytracing(unsigned char* ptr, sphere* s) 
__global__ void raytracing(unsigned
char* ptr)
}ptr[offset * 3 + 2] = (int)(r * 255);
ptr[offset * 3 + 1] = (int)(g * 255);
ptr[offset * 3 + 0] = (int)(b * 255);
}int main(int argc, char* argv)
//  cudamemcpy(s, temps, sizeof(sphere)*spheres, cudamemcpyhosttodevice);  
// 將球體引數copy進常量記憶體
cudamemcpytosymbol(s, temps, sizeof(sphere)*spheres);
free(temps);
dim3 grids(dim / 16, dim / 16);
dim3 threads(16, 16);
//  raytracing<<>>(devbitmap, s);    
raytracing << > > (devbitmap);
cudamemcpy(bitmap.data, devbitmap, 3 * bitmap.rows*bitmap.cols, cudamemcpydevicetohost);
cudaeventrecord(stop, 0);
cudaeventsynchronize(stop);
float elapsedtime;
cudaeventelapsedtime(&elapsedtime, start, stop);
printf("processing time: %3.1f ms\n", elapsedtime);
imshow("ray tracing", bitmap);
waitkey();
cudafree(devbitmap);
//  cudafree(s);    
return
0;}

實驗效果如下圖：

7.3 使用事件來測量效能

為了直觀地看到常量記憶體帶來的增益，我們需要測量程式執行的時間。

以往的話我們大多採用cpu或者作業系統中的某個計時器，但是這很容易帶來各種延遲(包括作業系統執行緒排程、高精度cpu計時器可用性等)。特別地，核函式與cpu程式是非同步執行的，這更易帶來意想不到的延遲。當然，針對這個問題，我們可以使用cudathreadsynchronize()函式進行同步然後再利用cpu計時。

除了採用cpu主機端計時之外，更準確的方法應該是利用cuda的事件api。

計時模板如下：

cudaevent_t start, stop;
float time = 0.f;
cudaeventcreate(&start);
cudaeventcreate(&stop);
cudaeventrecord(start, 0);
/*****
*****
*****
*****
*****
*****
*****
******
*****
*****
* 需要計時的**部分*
*****
*****
********
*****
*****
*****
*****
*****
*****
*******/ 
cudaeventrecord(stop, 0);
cudaeventsynchronize(stop);
cudaeventelapsedtime(&ime, start, stop);
std::cout << time
<< std::endl;
cudaeventdestroy(start);
cudaeventdestroy(stop);

參考料

GPU程式設計自學7 常量記憶體與事件

GPU程式設計自學3 CUDA程式初探

GPU程式設計自學4 CUDA核函式執行引數

GPU程式設計入門（7 效果（effect）基礎 2

GPU程式設計自學7 常量記憶體與事件

GPU程式設計自學3 CUDA程式初探

GPU程式設計自學4 CUDA核函式執行引數

GPU程式設計入門（7 效果（effect）基礎 2

相關推薦