CUDA學習(三十五)

2021-09-21 12:32:52 字數 2365 閱讀 7957

__constant__:

__constant__記憶體空間說明符(可選)與__device__一起使用,宣告乙個變數:

__shared__:

__shared__記憶體空間說明符,可以與__device__一起使用,

宣告乙個變數:

有塊的生命週期,

每塊有乙個獨特的物件,

只能從塊中的所有執行緒訪問

將共享記憶體中的變數宣告為外部陣列(如

extern __shared__ float shared;

陣列的大小在啟動時確定(請參閱執行配置)。 所有以這種方式宣告的變數,從記憶體中的相同位址開始,以便陣列中變數的布局必須通過偏移量進行顯式管理。 例如,如果有人想要相當於:

short array0[128];

float array1[64];

int array2[256];

在動態分配的共享記憶體中,可以通過以下方式宣告和初始化陣列:

extern __shared__ float array;

__device__ void func() // __device__ or __global__ function

請注意,指標需要與它們指向的型別對齊,因此下面的**不起作用,因為array1未對齊到4個位元組。

extern __shared__ float array;

__device__ void func() // __device__ or __global__ function

表3列出了內建向量型別的對齊要求

__managed__記憶體空間說明符(可選與__device__一起使用)宣告乙個變數:

__restrict__:

nvcc通過__restrict__關鍵字支援限制指標;

在c99中引入了受限制的指標以緩解存在於c型語言中的混疊問題,並且抑制了從**重新排序到常見子表示式消除的所有型別的優化

這是乙個受到別名問題影響的例子,使用受限制的指標可以幫助編譯器減少指令的數量:

void foo(const float* a,

const float* b,

float* c)

在c型語言中,指標a,b和c可能是別名,所以任何通過c寫入都可以修改a或b的元素。 這意味著為了保證函式的正確性,編譯器不能將[0]和b [0]載入到暫存器中,將它們相乘,並將結果儲存到c [0]和c [1]中,因為結果會與 抽象執行模型,例如,如果[0]與c [0]確實位於相同的位置。 所以編譯器不能利用公共的子表示式。 同樣,編譯器不能只將c [4]的計算重新排列到c [0]和c [1]的計算鄰近位置,因為前面寫入c [3]可能會將輸入更改為計算c [4]。

通過製作a,b和c限制指標,程式設計師向編譯器斷言指標實際上不是別名,在這種情況下意味著通過c寫入不會覆蓋a或b的元素。 這改變了函式原型如下:

void foo(const float* __restrict__ a,

const float* __restrict__ b,

float* __restrict__ c);

請注意,所有指標引數都需要被編譯器優化器限制,以獲得任何好處。 通過新增__restrict__關鍵字,編譯器現在可以重新排序並隨意執行常見的子表示式刪除,同時保留與抽象執行模型相同的功能:

void foo(const float* __restrict__ a,

const float* __restrict__ b,

float* __restrict__ c)

這裡的效果是減少了記憶體訪問次數,減少了計算次數。 這是由於「快取」的負載和常見的子表示式的登記壓力的增加而平衡的。

由於暫存器壓力在許多cuda**中是乙個關鍵問題,由於佔用率降低,使用受限制的指標可能會對cuda**產生負面的效能影響。

CUDA學習(三十五)

建議和最佳做法 整體效能優化策略 效能優化圍繞三個基本策略展開 最大限度地平行執行 優化記憶體使用量以實現最大記憶體頻寬 優化指令使用率以實現最大指令吞吐量 最大化並行執行從構建演算法開始,盡可能多地暴露資料並行。一旦演算法的並行性暴露出來,它就需要盡可能有效地對映到硬體。這是通過仔細選擇每個核心啟...

出差(三十五)

今天早晨到會議室後便買好了明天回公司的車票,長達乙個多月的出差生活即將告一段落,利用假期回去休息幾天,換換心情。經過這乙個多月的封閉式開發,對專案,對團隊有了進一步了解,感受到了創業公司的不易,也再次體驗了三點一線的生活方式,已經遠超996的工作模式,即使這樣依舊離目標有很大距離,主要原因還是缺少乙...

隨筆三十五 迴圈

1 do while 迴圈 do while 語句的通用形式如下 do body statement while test expr do while 的通用形式可以翻譯成如下所示的條件和 goto 語句 loop body statement t test expr if t goto loop ...