1. Introduction
近日,在写一些microbenchmark分析cuda程序访存问题时,发现了一个有趣的问题。目前尚未找到合理的解释,先记录下来以待后续分析。
实验平台为:NVIDIA GTX950,sm5.0,maxwell架构。
2. Global Memory
A memory "request" is an instruction which accesses memory, and a "transaction" is the movement of a unit of data between two regions of memory. Efficient access patterns minimize the number of transactions incurred by a request.
一般来说Global Memory的请求都要走Cache,部分架构的默认配置默认不仅走L2 cache,还要走L1 cache。如果L1、L2都被使用,那么memory transaction的单位是128bytes,如果仅使用L2 cache,则其基本单位是32bytes。如官方文档中的例子:
2.1 An example
以cuda sdk example中的vector_add代码为例:
__global__ void
vectorAdd(const float *A, const float *B, float *C, int numElements)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < numElements)
{
C[i] = A[i] + B[i];
}
}
设置numElements = 1
,则程序只有一个线程真正做事,计算C[0] = A[0] + B[0]
,一共有2次memory read和1次memory write请求。
设置numElements = 2
,则程序有两个线程真正做事,一共有4次memory read和2次memory write请求。
2.2 CUDA 9.0
设置numElements = 1
,使用nvprof --metrics gld_transactions,gst_transactions_per_request,gld_transactions_per_request,gst_transactions ./main
获得的profile结果如下表。
param | 解释 | 平均 |
---|---|---|
gld_transactions_per_request | Global Load Transactions Per Request | 1.000000 |
gst_transactions_per_request | Global Store Transactions Per Request | 1.000000 |
gld_transactions | Global Load Transactions | 2 |
gst_transactions | Global Store Transactions | 1 |
设置numElements = 2
:
param | 解释 | 平均 |
---|---|---|
gld_transactions_per_request | Global Load Transactions Per Request | 1.000000 |
gst_transactions_per_request | Global Store Transactions Per Request | 1.000000 |
gld_transactions | Global Load Transactions | 4 |
gst_transactions | Global Store Transactions | 2 |
ldst_issued | Issued Load/Store Instructions | 19 |
tex_cache_transactions | Unified Cache Transactions | 2 |
实际上,4次memory load在真正进行内存操作时,可以进行内存请求合并(memory coalescing),将A[0]和A[1]的读取合并为一个transaction。
2.3 CUDA 10
使用cuda10,cuda9.2编译,设置numElements = 2
。
param | 解释 | 平均 |
---|---|---|
gld_transactions | Global Load Transactions | 10 |
gld_transactions_per_request | Global Load Transactions Per Request | 5.000000 |
gst_transactions | Global Store Transactions | 1 |
global_load_requests | Total number of global load requests from Multiprocessor | 2 |
inst_executed_global_loads | Warp level instructions for global loads | 2 |
ldst_issued | Issued Load/Store Instructions | 5 |
tex_cache_transactions | Unified Cache Transactions | 2 |
3. 分析
3.1 问题1 L1 cache问题
sm5.0可以使用-Xptxas -dlcm=ca
和cg
来决定是否让global memory的请求被L1 cache,而且也提到sm5.0默认global memory数据只被L2 cache。但是从§2的nvprof结果中可以看到,tex_cache_transactions一直为2,不论使用ca还是cg,这个有点奇怪,想不明白为什么。
3.2 问题2
2.3中的global_load_requests看起来已经经过了内存请求的合并,但是请求的是A[0], A[1], B[0], B[1],一共两个8bytes。而cuda programing guide中表示sm5.0 global memory默认仅被L2 cached,因此一个transaction为32bytes,足够cover 8bytes的请求。所以按理说,两个内存请求,用两个transaction就可以完成,但是不知道为什么cuda10和cuda9.2的global load transactions是10,per request是5。
而且,cuda9.0的profile结果中,gld_transactions = 4,而不是上面我们分析的2。
后续如果找到合理的解释,再更新文章。
Reference
CUDA之Global memory合并访问Coalesced详解
L2 Hit Rate(Texture Reads) becomes 100% when modifying memory never used
Comments