3

最近我正在研究CUDA。我想知道 CUDA 内存访问时间。

在 CUDA Programming Guide 中写了内存访问时间：

  • 全局内存访问时间为 400 ~ 600 Cycle
  • 共享内存（L1 Cache）访问时间为 20 ~ 40 Cycle

我认为 Cycle 与时钟相同。它是否正确 ？如果那是正确的，那么我检查了内存访问时间。主机是固定的，但内核代码有 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 );

结果

全局内存访问#1：882
全局内存访问#2：312
全局内存访问#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 );

结果

全局内存访问#1：872
全局内存访问#2：776
全局内存访问#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 );

结果

全局内存访问#1：168
全局内存访问#2：168
全局内存访问#3：168

我不明白这个结果

Dev_In[50000]、Dev_In[99999]没有缓存，但是访问时间非常快！！只是，我使用了 1 个变量....

所以..我的问题是gpu周期== gpu时钟？

在 result1、result2、result3 中，为什么 result3 中的内存访问时间非常快？

4

1 回答 1

1

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

于 2012-09-19T10:57:45.537 回答