> For the complete documentation index, see [llms.txt](https://legacy.cookielau.com/llms.txt). Markdown versions of documentation pages are available by appending `.md` to page URLs; this page is available as [Markdown](https://legacy.cookielau.com/archives/10-hardware/0-gpu/008-memoryclassification.md).

# Memory Classification

Memory Classification

2020 Jan 29th CookieLau

## Preface

![](/files/-M78qX1N9cfQ8psLFwpi)

## 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 ![](/files/-M78qX1bpmvuNo2zd7L7)

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

![](/files/-M78qX1k4Mm4lDQPocFk)

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

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

```c
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

1.