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=cacg来决定是否让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

sm3.0的global memory

CUDA之Global memory合并访问Coalesced详解

Memory Transactions

SM5.0的global memory

L2 Hit Rate(Texture Reads) becomes 100% when modifying memory never used


文章版权归 FindHao 所有丨本站默认采用CC-BY-NC-SA 4.0协议进行授权|
转载必须包含本声明,并以超链接形式注明作者 FindHao 和本文原始地址:
https://findhao.net/academic/2564.html

Comments