关闭

cudaMemcpyToSymbol 和cudaMemcpy

68人阅读 评论(0) 收藏 举报
分类:

这是今天在群里有人问了这么一个问题 
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的准确率,应该是指性能模拟。计算结果准确性的话,在研究方面不突出……

0
0

查看评论
* 以上用户言论只代表其个人观点,不代表CSDN网站的观点或立场
    个人资料
    • 访问:62571次
    • 积分:2846
    • 等级:
    • 排名:第12667名
    • 原创:197篇
    • 转载:233篇
    • 译文:1篇
    • 评论:10条
    坐标
    天马脚下,岳麓山前
    文章分类
    最新评论