c++ 全局和共享内存访问时间

rqdpfwrv  于 2023-05-20  发布在  其他
关注(0)|答案(1)|浏览(106)

最近我在学习CUDA。我想知道CUDA内存访问时间。
在《CUDA编程指南》中写的内存访问次数:

  • 全局存储器访问时间为400 ~ 600周期
  • 共享内存(L1 Cache)访问时间为20 ~ 40周期

我认为周期和时钟是一样的。是这样吗?如果这是正确的,那么我检查了内存访问时间。主机是固定的,但内核代码有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时钟?
为什么结果3中的内存访问时间非常快?

vd8tlhqk

vd8tlhqk1#

由于@phoad所述的原因,您的评估无效。在内存访问之后和时钟停止之前,您应该重用内存读取值,以使指令依赖于未完成的负载。否则,GPU一个接一个地发出独立指令,并且在时钟开始和加载之后立即执行时钟结束。我建议你试试亨利Wong在here准备的微基准测试套装。使用这个套件,您可以检索各种微体系结构的详细信息,包括内存访问延迟。如果您只需要内存延迟,那么可以尝试Sylvain Collange开发的CUDA latency

相关问题