Memory Classification

Memory Classification

2020 Jan 29th CookieLau

Preface

Register

Local Memory

Shared Memory

详见 014 Using Shared Memory

Global Memory

Global Memory 讲究的是 coalesce ，也就是合并存取，即在 global memory 上面访存每次都是取连续的 32-bytes 的倍数的段，所以当可以合并存取的时候就合并存取，有利于提高带宽，举个例子： 一个 warp 中的 threads 访存连续的 4-bytes 大小的数据，则 global memory 上存取的可能是如下： 从 96 到 224 bytes 一共 32*4=128 bytes 都被取到，一共取了 4个 32-bytes，是 coalesced access

而如果比如说一个 warp 中的 threads 取连续的但不是 32-bytes 的倍数的大小的数据时候，可能多出来的不够 32-bytes 的部分 仍要 从 global memory 处拿 32-bytes，这样就有一部分的浪费，导致带宽的消耗。

Cuda 自身的 API 是很注重对齐的，比如 cudaMalloc 分配的内存都是对齐 256-bytes，所以我们在访存的时候也要注意这个特点，比如将 block thread 的size分配的合理一些，warp 的倍数，32的倍数之类的。

比如比较推荐的一种方法就是：

int deviceId;
int numberOfSMs;

cudaGetDevice(&deviceId);
cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);

size_t threadsPerBlock;
size_t numberOfBlocks;

threadsPerBlock = 256;
numberOfBlocks = 32 * numberOfSMs;

kernel<<<numberOfBlocks, threadsPerBlock>>>(...)

其实大小没有很大的影响，32的倍数即可。

Constant Memory

Texture Memory

Reference