最近在研究CUDA。我想了解 CUDA 内存访问时间。
在《CUDA Programming Guide》中写的内存访问次数:
- 全局内存访问时间为400 ~ 600 Cycle
- 共享内存(L1 Cache)存取时间为20~40 Cycle
我认为 Cycle 与时钟相同。这样对吗 ?如果那是正确的,那么我检查了内存访问时间。主机是固定的,但内核代码有 3 个版本。这是我的代码:
host Code
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();
kernel version 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] 被缓存了..
kernel version 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 迟到了...
kernel version 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的内存访问时间很快?
最佳答案
由于@phoad 所述的原因,您的评估无效。在内存访问之后和时钟停止之前,您应该重新使用内存读取值以使指令依赖于未完成的负载。否则,GPU 会一个接一个地发出独立的指令,并且在时钟启动和加载后立即执行时钟结束。我建议你试试 Henry Wong 在 here 准备的微基准测试套装。 .使用此套装,您可以检索各种微体系结构详细信息,包括内存访问延迟。如果你只需要内存延迟,尝试CUDA latency会更容易由 Sylvain Collange 开发。
关于cuda 全局和共享内存访问时间,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/12480696/