cudaMemcpyToSymbol 和cudaMemcpy

这是今天在群里有人问了这么一个问题 
cudaMemcpyToSymbol可以将数据从host拷贝到global，cudaMemcpy也是从host到>global，这种情况下二个函数有什么区别吗？

和各位大佬讨论一下后，和大家分享一下~

首先，学到了cudaMemcpyToSymbol竟然还有将数据从host拷贝到global的功能，以前只用过这个函数拷贝到constant memory。拷贝方式的不同是由目的内存申请的方式决定的。

申请的是device内存，cudaMemcpyToSymbol拷贝就是从host拷贝到global memory。

申请的是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的准确率，应该是指性能模拟。计算结果准确性的话，在研究方面不突出……