版主你好，在一台主机情况下，多个gpu在运行过程中需要交换某些迭代数据，此时用主机内存作为媒介是不是比较慢？有什么方法达到更快么？网上有人说用p2p或MPI。恳请版主指点迷津。。。。
楼主您好，
如果需要普适的回答，那么您只能通过内存交换数据。这是唯一适用所有计算能力的卡的方式。
如果您没有老的1.x的卡，那么您可以使用P2P copy (cudaMemcpyPeer), 这个只要是计算能力>=2.0的都能支持，可以直接在您的单CPU主机上的多张卡之间通过PCI-E总线交换数据，不走内存。
以及，如果您的>=2.0计算能力的卡，同时还是Tesla品牌的，您可以可以考虑直接使用P2P access, 此时一个卡上的kernel将有能力访问另外一张卡上的数据，无需单独手工复制。对于小数据量或者不方便进行复制的场合，用起来相当方便。建议您考虑选购。
感谢您的来访。
玫瑰幻想版主，
您好,看了你的解答我很开心，因为我们实验室准备购买的就是Tesla品牌的K20。但是我还有两个问题：
一，P2P access如何快速上手，可否指明下道路？
二，一个卡上的kernel访问另外一张卡上的数据时，速度如何，会不会时间开销也很大？因为有可能1卡还没来得及访问2卡，2卡的数据就被更新了。
有劳版主了！
(1)举个例子，如果楼主您有2张卡，分别是卡0和卡1,那么:
在您的程序开头：
cudaSetDevice(0);
cudaDeviceEnablePeerAccess(1,0); //当前设备（卡0）可以访问卡1了。
然后：
cudaSetDevice(1);
cudaDeviceEnablePeerAccess(0,0); //当前设备（卡1）可以访问卡0了。
然后您可以在任何一个卡上cudaMalloc分配显存，然后分配出来的显存可以在任何一个卡上的kernel中使用，例如：
cudaSetDevice(0);
int *p = NULL;
cudaMalloc(&p, …); //p是设备0上分配的一段缓冲区
cudaSetDevice(1);
kernel<<<>>>(p,…); //但是却可以在卡1上的kernel中使用此段显存
（反过来也是的，只要你开头加上去刚才的那些语句）
（2）这个实际上不是一个问题，也不是硬件的责任，实际上这和您在1张卡上同时运行2个kernel一个道理。
要如果确保一个kernel没更新完数据前，另外一个kernel不要启动，需要您自行进行一定的同步处理，在正确的时刻进行访问。或者您可以选择不要随意更新共享的数据（将共享的数据只读）。
（例如：您可以使用cudaEventSynchronize进行同步（event同步可以跨卡的））
（例如：您可以将特定的缓冲区不要就地更新，而将结果保存到另外一个缓冲区，这样对源缓冲区只读）
感谢您的深夜来访。
如果贵试验室需要购买K20, 建议您（或者您推荐贵实验室的购买负责人）联系：
吉浦迅科技，顾燕，电话：18915751925
吉浦迅科技是NVIDIA首席合作伙伴，提供NVIDIA显卡试用，显卡销售，售后技术支持等服务。
顾燕是我的朋友。
玫瑰幻想版主，
很感激您耐心又周详的解答，我会按照您的建议继续学习的，至于购买设备问题，这个我实在无能为力，但期待以后会有所合作，非常感谢您的美意！
感谢您的来访，期待您的下次光临。