CUDA 全局和共享内存访问时间

Question

最近我在学习CUDA。我想了解一下CUDA内存访问时间。

CUDA编程指南中写有内存访问时间：

• 全局内存访问时间为400~600个周期
• 共享内存（L1缓存）访问时间为20~40个周期

我认为“Cycle”和“clock”是相同的。这正确吗？如果是正确的，那么我就研究一下内存访问时间。主机是固定的，但内核代码有3个版本。这是我的代码：

---

主机代码：

float* H1  = (float*)malloc(sizeof(float)*100000);
float* D1;
    
for( int i = 0 ; i < 100000 ; i++ ){
    H1[i]  = i;
}
        
cudaMalloc( (void**)&D1,   sizeof(float)*100000);
cudaMemcpy( D1, H1,    sizeof(float)*100000, cudaMemcpyHostToDevice );
            
    
cudaPrintfInit();
test<<<1,1>>>( D1 );
cudaPrintfDisplay(stdout, true);

cudaPrintfEnd();

---

内核版本 1：

float Global1;
float Global2;
float Global3;

int Clock;

Clock = clock();
Global1 = Dev_In1[1];
Clock = clock() - Clock;
cuPrintf("Global Memory Access #1 : %d\n", Clock );

Clock = clock();
Global2 = Dev_In1[2];
Clock = clock() - Clock;
cuPrintf("Global Memory Access #2 : %d\n", Clock );

Clock = clock();
Global3 = Dev_In1[3];
Clock = clock() - Clock;
cuPrintf("Global Memory Access #3 : %d\n", Clock );

这是结果

Global Memory Access #1 : 882       
Global Memory Access #2 : 312       
Global Memory Access #3 : 312       

我认为第一次访问没有缓存，所以需要800个周期，但第二次和第三次访问只需要312个周期，因为Dev_In[2]和Dev_In[3]已经被缓存了...

---

内核版本 2：

int Global1, Global2, Global3;              
int Clock;              
            
Clock = clock();                
Global1 = Dev_In1[1];               
Clock = clock() - Clock;                
cuPrintf("Global Memory Access #1 : %d\n", Clock );             
            
Clock = clock();                
Global2 = Dev_In1[50000];               
Clock = clock() - Clock;                
cuPrintf("Global Memory Access #2 : %d\n", Clock );             
            
Clock = clock();                
Global3 = Dev_In1[99999];               
Clock = clock() - Clock;                
cuPrintf("Global Memory Access #3 : %d\n", Clock );             

这是结果

Global Memory Access #1 : 872       
Global Memory Access #2 : 776       
Global Memory Access #3 : 782       

我认为在第一次访问时，Dev_In1[50000]和Dev_In2[99999]没有被缓存

所以... #1，#2，#3都晚了...

---

内核版本 3：

int Global1, Global2, Global3;                  
int Clock;                  
                
Clock = clock();                    
Global1 = Dev_In1[1];                   
Clock = clock() - Clock;                    
cuPrintf("Global Memory Access #1 : %d\n", Clock );                 
                
Clock = clock();                    
Global1 = Dev_In1[50000];                   
Clock = clock() - Clock;                    
cuPrintf("Global Memory Access #2 : %d\n", Clock );                 
                
Clock = clock();                    
Global1 = Dev_In1[99999];                   
Clock = clock() - Clock;                    
cuPrintf("Global Memory Access #3 : %d\n", Clock );                 

结果：

Global Memory Access #1 : 168       
Global Memory Access #2 : 168       
Global Memory Access #3 : 168       

我不明白这个结果。

Dev_In[50000] 和 Dev_In[99999] 没有被缓存，但访问速度非常快！只是，我使用了一个变量...

那么...我的问题是GPU周期是否等于GPU时钟？

为什么第三个结果的内存访问时间非常快？

- 오승택

---

你的编译器可能会对代码进行优化。也就是说，因为你没有在任何地方使用Global1，你的编译器可能不会将其全部读取。在从内存中读取下一个值之前，请尝试使用Global1。 - phoad

clock()函数返回的值应该与设备属性CU_DEVICE_ATTRIBUTE_CLOCK_RATE指定的值相同。clock()函数返回的值是以周期为单位的，它不是固定频率的。在执行内核时，2.x及更早版本的设备上频率是固定的。在3.x设备上，时钟速率会随着电源管理和热管理而变化。 - Greg Smith

你是在同一个CUDA上下文中同时执行这三个内核，还是启动进程3次？如果是前者，则使用相同的设备内存运行kernel2然后运行kernel3可能会导致值被缓存在L2中，这将解释显著较低的经过时间。使用clock()时，应验证SASS（汇编代码，而不是PTX）具有正确的指令序列。编译器通常会移动或消除对clock()的调用。我建议在cuPrintf之后插入threadfence()以刷新LSU单元。 - Greg Smith

Answer 1

由于@phoad所述的原因，您的评估无效。在内存访问之后和时钟停止之前，您应该重用内存读取值以使指令依赖于未完成的负载。否则，GPU会独立执行一系列指令，并且时钟结束会立即在时钟开始和负载之后执行。我建议您尝试Henry Wong在这里准备的微基准套件。使用此套件，您可以检索各种微架构细节，包括内存访问延迟。如果您只需要内存延迟，则更容易尝试由Sylvain Collange开发的CUDA延迟。

- lashgar