網頁

2019年2月6日 星期三

5.3.2. Device Memory Accesses (To be continuous)

記憶讀取可能會需要傳輸很多次,這取決於warp內的thread它們如何執行記憶體addresses分佈。例如global memory的通則是addresses分散越多,throughput越少。


Global Memory
global memory到device memory是經由32,64或128bytes memory transactions進行傳輸。例如,如果為每個thread 4 bytes存取,throughput分成八次的話,可以得到32 bytes的memory transactions。



Size and Alignment Requirement
若且唯若資料型態大小為1, 2, 4, 8,或16 bytes,而且資料自然地被對齊(address大小為前述的倍數),則對停留在global memory的數據其任何存取(經由variable或pointer)都編譯為單一global memory指令。

對於structure,編譯器可以使用對齊specifiers __align __(8)或 __align __(16)強制執行大小和對齊要求,例如

struct __align__(8) {
    float x;
    float y; 
}; 


struct __align__(16) {
    float x;
    float y;
    float z; 
}; 

停留在global memory或者是由驅動或runtime API 中的一個memory allocation routines返回的variable其任何address始終對齊至少256個bytes。

讀取非自然對齊的8 byte或16 bytes words會產生不正確的結果,因此必須特別注意保持這些類型的任何值或陣列的起始address的對齊。

以下這個典型情況是當使用一些自定義global memory allocation時,多組陣列的allocation(多次呼叫cudaMalloc()或cuMemAlloc())被單獨大的block memmory分割成多組陣列所取代 ,在這種情況下,可能容易忽視每個的起始地址都已經偏離原本大的block的起始地址。













沒有留言:

張貼留言