cudaMemcpyToSymbol 和cudaMemcpy

转载 2017年01月03日 15:34:23

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

拷贝global memory,cudaMemcpyToSymbol 和cudaMemcpy函数是否有区别

这是今天在群里有人问了这么一个问题 cudaMemcpyToSymbol可以将数据从host拷贝到global,cudaMemcpy也是从host到>global,这种情况下二个函数有什么区别吗...
  • litdaguang
  • litdaguang
  • 2015年04月14日 19:56
  • 2649

CUDA学习笔记九

 Memory kernel性能高低是不能单纯的从warp的执行上来解释的。比如之前博文涉及到的,将block的维度设置为warp大小的一半会导致load efficiency降低,这个问题无...
  • langb2014
  • langb2014
  • 2016年05月08日 23:47
  • 6248

cudaMemcpyToSymbol

 __constant__ int maxI;int dummy = 1;CUDA_SAFE_CALL( cudaMemcpyToSymbol(maxI, &dummy, 1 * sizeof(int...
  • dvchn
  • dvchn
  • 2008年03月02日 14:34
  • 1605

cudaMemcpyToSymbol()的invalid device symbol问题

cudaMemcpyToSymbol()的invalid device symbol问题解决方法,最后解决方案是去掉函数第一个参数的&符号。虽然最后知道了方法很简单,但寻找解决方案的过程很复杂,因而做...
  • sysysty
  • sysysty
  • 2017年04月13日 09:46
  • 754

cudaMemcpy与cudaMemcpyAsync的区别

简单可以理解为:cudaMemcpy是同步的,而cudaMemcpyAsync是异步的。具体理解需要弄清以下概念: 1.CUDA Streams 在cuda中一个Stream是由主机代码发布的一系...
  • Bruce_0712
  • Bruce_0712
  • 2017年06月19日 21:06
  • 664

cudaMemcpy与kernel

Memory操作 cuda程序将系统区分成host和device,二者有各自的memory。kernel可以操作device memory,为了能很好的控制device端内存,CUDA提供了几个...
  • u014800094
  • u014800094
  • 2016年10月12日 18:19
  • 818

CUDA全局变量(__device__)的初始化与使用:cudaMemoryToSymbol、cudaMemoryFromSymbol、cudaGetSymbolAddress

在cuda中在设备(device)中声明一个全局变量用__device__关键字修饰: __device__ float devData; 初始化为: float value = 3.14f; cud...
  • Rong_Toa
  • Rong_Toa
  • 2017年11月29日 14:24
  • 120

cudaMemcpyToSymbol 和cudaMemcpy

这是今天在群里有人问了这么一个问题  cudaMemcpyToSymbol可以将数据从host拷贝到global,cudaMemcpy也是从host到>global,这种情况下二个函数有什么区别吗?...
  • u014800094
  • u014800094
  • 2017年01月03日 15:34
  • 231

CUDA中的常量内存__constant__

GPU包含数百个数学计算单元,具有强大的处理运算能力,可以强大到计算速率高于输入数据的速率,即充分利用带宽,满负荷向GPU传输数据还不够它计算的。CUDA C除全局内存和共享内存外,还支持常量内存,常...
  • dcrmg
  • dcrmg
  • 2017年02月06日 19:21
  • 1871

CPU 与 GPU 之间数据转换 cudaMemcpy

//显存上分配空间 CUDA_SAFE_CALL(cudaMalloc((void**)&Dst_d,sizeof(float3) * totalPNum)); //显存上传输数据 CUDA_S...
  • xiaoheibaqi
  • xiaoheibaqi
  • 2015年03月01日 10:22
  • 498
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:cudaMemcpyToSymbol 和cudaMemcpy
举报原因:
原因补充:

(最多只允许输入30个字)