CUDA Memory Model

對于程序員來說，memory可以分為下面兩類：

• Programmable：我們可以靈活操作的部分。
• Non-programmable：不能操作，由一套自動機制來達到很好的性能。

在CPU的存儲結構中，L1和L2 cache都是non-programmable的。對于CUDA來說，programmable的類型很豐富：

• Registers
• Shared memory
• Local memory
• Constant memory
• Texture memory
• Global memory

下圖展示了memory的結構，他們各自都有不用的空間、生命期和cache。

Global Memory

global Memory是空間最大，latency最高，GPU最基礎的memory。“global”指明了其生命周期。任意SM都可以在整個程序的生命期中獲取其狀態。global中的變量既可以是靜態也可以是動態聲明。可以使用__device__修飾符來限定其屬性。global memory的分配就是之前頻繁使用的cudaMalloc，釋放使用cudaFree。global memory駐留在devicememory，可以通過32-byte、64-byte或者128-byte三種格式傳輸。這些memory transaction必須是對齊的，也就是說首地址必須是32、64或者128的倍數。優化memory transaction對于性能提升至關重要。當warp執行memory load/store時，需要的transaction數量依賴于下面兩個因素：

Distribution of memory address across the thread of that warp 就是前文的連續

Alignment of memory address per transaction 對齊

一般來說，所需求的transaction越多，潛在的不必要數據傳輸就越多，從而導致throughput efficiency降低。

對于一個既定的warp memory請求，transaction的數量和throughput efficiency是由CC版本決定的。對于CC1.0和1.1來說，對于global memory的獲取是非常嚴格的。而1.1以上，由于cache的存在，獲取要輕松的多。

下面代碼是通過可以通過32-byte、64-byte或者128-byte三種格式傳輸，優化傳輸效率：

template <typename T> __device__ inline uint32_t pack_uint8x4(T x, T y, T z, T w){uchar4 uint8x4;uint8x4.x = static_cast<uint8_t>(x);uint8x4.y = static_cast<uint8_t>(y);uint8x4.z = static_cast<uint8_t>(z);uint8x4.w = static_cast<uint8_t>(w);return load_as<uint32_t>(&uint8x4); }template <unsigned int N> __device__ inline void store_uint8_vector(uint8_t *dest, const uint32_t *ptr);template <> __device__ inline void store_uint8_vector<1u>(uint8_t *dest, const uint32_t *ptr){dest[0] = static_cast<uint8_t>(ptr[0]); }template <> __device__ inline void store_uint8_vector<2u>(uint8_t *dest, const uint32_t *ptr){uchar2 uint8x2;uint8x2.x = static_cast<uint8_t>(ptr[0]);uint8x2.y = static_cast<uint8_t>(ptr[0]);store_as<uchar2>(dest, uint8x2); }template <> __device__ inline void store_uint8_vector<4u>(uint8_t *dest, const uint32_t *ptr){store_as<uint32_t>(dest, pack_uint8x4(ptr[0], ptr[1], ptr[2], ptr[3])); }template <> __device__ inline void store_uint8_vector<8u>(uint8_t *dest, const uint32_t *ptr){uint2 uint32x2;uint32x2.x = pack_uint8x4(ptr[0], ptr[1], ptr[2], ptr[3]);uint32x2.y = pack_uint8x4(ptr[4], ptr[5], ptr[6], ptr[7]);store_as<uint2>(dest, uint32x2); }template <> __device__ inline void store_uint8_vector<16u>(uint8_t *dest, const uint32_t *ptr){uint4 uint32x4;uint32x4.x = pack_uint8x4(ptr[ 0], ptr[ 1], ptr[ 2], ptr[ 3]);uint32x4.y = pack_uint8x4(ptr[ 4], ptr[ 5], ptr[ 6], ptr[ 7]);uint32x4.z = pack_uint8x4(ptr[ 8], ptr[ 9], ptr[10], ptr[11]);uint32x4.w = pack_uint8x4(ptr[12], ptr[13], ptr[14], ptr[15]);store_as<uint4>(dest, uint32x4); }

?例子代碼見：

https://github.com/Alexjqw/cuda-learning

?

參考：https://www.cnblogs.com/1024incn/p/4564726.html

?

總結

以上是生活随笔為你收集整理的cuda的global memory介绍的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。