# How\_to\_Access\_Global\_Memory\_Efficiently
How to Access Global Memory Efficiently in CUDA C/C++ Kernels
2020 Jan 28th CookieLau
Source:
\[toc]
## testCode
分别用一下代码测试 offset 访问和 stride 访问的间隔对带宽的影响：
```cpp
__global__ void offset(T* a, int s)
{
int i = blockDim.x * blockIdx.x + threadIdx.x + s;
a[i] = a[i] + 1;
}
__global__ void stride(T* a, int s)
{
int i = (blockDim.x * blockIdx.x + threadIdx.x) * s;
a[i] = a[i] + 1;
}
```
## Misaligned Data Access
C870: Compute Capability 1.0 C1060: Compute Capability 1.3 C2050: Compute Capability 2.0
在 device 中分配的数组都被 cuda Driver 按照 256字节 对齐，当访存 global Memory 的时候可以通过 32字节、64字节 或 128字节等分块进行数据交换。
对于C870这种 compute capability 在 1.0 及以下的 GPU，其 warp size 只有16，而且当发生 misaligned access 的时候，会对每个 misaligned 的 data 单独进行存取，所以会从原来的 16 thread 变成 16次的 32-bytes 访存。对于存取 float 数据而言，每次取 32-bytes 的数据中只有 4-bytes 是有效的，所以带宽减少到原来(offset=0) 的 1/8.
对于C1060这种 compute capability 较好的，对带宽的影响没有 C870 那么严重，只要 misaligned 落在访存的 segment，如 32,64,128-bytes 只会降低部分的性能。
对于C2050这种 compute capability 大于等于 2.0 的设备而言，其每个 multiprocessor 都配有一个 128-bytes 的 L1 cache，所以offset的改变几乎不会对带宽产生影响。
## Stride Memory Access
对于C870完全不能处理misaligned，只能处理linear+aligned的架构来说，除了stride=1的情况，都出现了 7/8 的落差，变为正常情况下的 1/8。
对于CC(Compute Capability>1.0) 的来说，都能处理部分 misaligned，所以带宽曲线是 smoothly 下降，但对于相隔很远的访问也无能为力。
但是我们又时常需要进行 stride 访问，比如 grid-stride，则可以通过 **shared memory** 来解决这一问题。 Shared memory 是 on-chip 的、**被一个 thread block 中的所有 threads** 共享的一部分内存。
举一个例子：将 2D 的数组的数据加载到 shared memory 里面进行访存降低对带宽的伤害。在共享内存中是没有 stride access penalty 的。
---
# Agent Instructions: Querying This Documentation
If you need additional information that is not directly available in this page, you can query the documentation dynamically by asking a question.
Perform an HTTP GET request on the current page URL with the `ask` query parameter:
```
GET https://legacy.cookielau.com/archives/10-hardware/0-gpu/013-how_to_access_global_memory_efficiently.md?ask=
```
The question should be specific, self-contained, and written in natural language.
The response will contain a direct answer to the question and relevant excerpts and sources from the documentation.
Use this mechanism when the answer is not explicitly present in the current page, you need clarification or additional context, or you want to retrieve related documentation sections.