这是今天在群里有人问了这么一个问题
cudaMemcpyToSymbol
可以将数据从host拷贝到global,cudaMemcpy
也是从host到>global,这种情况下二个函数有什么区别吗?
和各位大佬讨论一下后,和大家分享一下~
首先,学到了cudaMemcpyToSymbol
竟然还有将数据从host拷贝到global的功能,以前只用过这个函数拷贝constant memory。拷贝方式的不同是由目的内存申请的方式决定的。申请的是device内存,cudaMemcpyToSymbol
拷贝就是从host拷贝到global。申请的是constant内存,cudaMemcpyToSymbol
拷贝就是从host拷贝到constant memory。
__device__ float g_damp_x[MAX_DIM];
__constant__ float c_damp_y[MAX_DIM];
然后大家开始各种猜测,性能方面的问题、设计方面的问题以及兼容性方面问题等等。下面是我的一些见解,很荣幸得到了各位大佬的赞同,跟大家分享一下。
我是从GPU性能仿真器GPGPU-Sim(http://www.gpgpu-sim.org)入手的,通过看Sim中这个两个函数部分的代码,可以看到这个两个函数的执行过程。
cudaMemcpy
函数的过程比较简单,只是做了一些安全判断后就像模拟global内存中根据“地址”将数据写到global内存中(GPGPU-Sim是通过物理内存来模拟GPU内的各种内存)。
而cudaMemcpyToSymbol
函数的过程实现的比较复杂,具体在cuda-sim/cuda-sim.cc
文件的1554行的gpgpu_ptx_sim_memcpy_symbol
函数中,在这个函数中,在根据目标地址在global memory和constant memory中进行搜索。代码如下:
if( g_globals.find(hostVar) != g_globals.end() ) {
found_sym = true;
sym_name = hostVar;
mem_region = global_space;
}
if( g_constants.find(hostVar) != g_constants.end() ) {
found_sym = true;
sym_name = hostVar;
mem_region = const_space;
}
然后再根据搜索结果将数据写到不同的内存中,所以根据以上观察的结果,cudaMemcpy在性能会更好一些。
PS:今天请教了实验室里的学长这个问题,问题大概是这样:
1.Sim的设计者是如何知道NVIDIA公司GPU产品的底层实现,这个技术细节不是商业机密么?
GPGPU-Sim是一个开放性的实验平台,它在仿真方面有很高的准确性,但是在具体细节方面,我们可以认为他是合理的,但不一定是绝对准确的。因为sim的设计中有NVIDIA的人,在文档中有提到过。而且Sim所能使用的cuda还是4.0,所以NVIDIA可能会通过这个公布一些老版本架构里的细节,用于科研研究。
2.如果手上没有GPU,可以装这个仿真器来代替GPU跑cuda程序吗?
这个仿真器确实可以很好模拟比较老的架构,但是因为是串行模拟并行,所以一个在真实GPU上几秒就跑完的程序,在性能模拟模式下,在Sim里大概需要几个小时甚至一天 。但是Sim的配置文件里有一个选项(pure functional simulation:-gpgpu_ptx_sim_mode),可以只功能模拟,不输出性能信息,那个还是比较快的。而且根据sim文档中描述的,在GT200架构上的准确率是98.37%,在Fermi上的准确率是97.35%,已经很高了。
3.准确率是指什么?计算结果的准确性还是性能模拟的吻合度?
IPC的准确率,应该是指性能模拟。计算结果准确性的话,在研究方面不突出……