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

Last updated

Was this helpful?