关于cudaMemcpy2D的问题，谢谢！

[guoboxue986235]

int src_width = 200;
int src_height = 80;
int dst_width = 100;
int dst_height = 80;

int *pSrc,*pDst;
cudaMalloc(void(**)&pSrc, src_width * src_heihgt * sizeof(int));

cudaMalloc(void(**)&pDst, dst_width * dst_heihgt * sizeof(int));

我想把src的前100列COPY到目的中，如何用cudaMemcpy2D,
谢谢，试了一天了，都没有搞定，十分感谢！

[无心人_过过小日子]

1)不要用cudaMalloc,用cudaMallocPitch.
int spitch, dpitch;

cudaMallocPitch(void(**)&pSrc, &spitch, width*sizeof(int), src_height);
cudaMallocPitch(void(**)&pDst, &dpitch, width*sizeof(int), dst_height);

2)
cudaMemcpy2D(pDst, dpitch, pSrc, spitch, width, height,cudaMemcpyDeviceToDevice);
复制前height行.

如果要复制列的话,要写个内核程序执行转置.还不如直接自己写个复制内核呐.

[guoboxue986235]

谢谢！自己写的COPY函数，当然是开很多线程跑
有没有cudaMemcpy快？

还有一个问题，想问您：如果在两个显卡间COPY数据，而不经过内存。我按照您的指点，看了下SDK中的那个多显卡的例子，好像没有我想要的东西？谢谢

[无心人_过过小日子]

1)写的好的话,该差不多的.
2)不行的.

[OpenHero]

倒是主内存上可以用SSE来做memcopy 可能会快一些
但是在 device到device还是得通过host做中转

[天下第一好大人]

假定需求统一，就是简单的按列拷贝，源和目的都在cpu端（如果有后续计算，那就不一定了）

1> memory 开成 pinned，然后用微软的memcpy试试(不要用stdlib的)。
2> cudaMemcpy2D，属性HostToHost
3> cudaMemcpy到显存，然后cudaMemcpy2D，属性DeviceToDevice，然后cudaMemcpy回内存
4> 自己写内核实现。

楼主试试哪个快。主要是你的数据量太小，所以我估计可能直接在cpu上做就行。

[OpenHero]

1.SSE说的是host段的，主内存的copy；这个是GPU无关
2.GPU上的内存到内存copy，针对一个GPU当然会很快，device到device是硬件加速的，可以看看内部的架构，并不走技术单元，直接copy；
3.多GPU之间的copy只能通过host中转；

[guoboxue986235]

在内存上通过SSE来作memcopy也没有显存到显存的速度快，这个我测试过，显存到显存COPY的时候到底是怎么COPY的，为什么会那么快的？

[guoboxue986235]

谢谢，受教了！

[Cyrosly]

不用内核，一次性COPY列的话（你的例子中其实相当于COPY SRC GMEM的一个子区域到DST GMEM)，用driver API是可以的,想起了opengl中的glTexCopySubImage2D,是类似的:

planeMem.srcMemoryType =CU_MEMORYTYPE_DEVICE;
planeMem.srcDevice     =pSrc;
planeMem.srcXInBytes   =0;                  //src dmem start in the first row
planeMem.srcPitch      =200*sizeof(float);  //src span between two rows
planeMem.srcY          =0;                  //the start col of copy form
planeMem.dstMemoryType =CU_MEMORYTYPE_DEVICE;
planeMem.dstDevice     =pDst;                
planeMem.dstXInBytes   =0;                  //dst start byte in the first row of dmem
planeMem.dstY          =0;                  //the start col of copy to
planeMem.dstPitch      =100*sizeof(float);  //span len when copyed form src
planeMem.WidthInBytes  =100*sizeof(float);  //dst width
planeMem.Height=80;                         //dst height
cuMemcpy2D(&planeMem);

RT API肯定也是可以的，但我自己用DA 用习惯了，自己试试吧

[guoboxue986235]

好的，我试下，谢谢！

[无心人_过过小日子]

应该还是gpu上的内核copy的.只是用更低的api(没开放)写的吧.
另外,gpu上带宽比主板上大很多(多的话20倍吧),copy快很正常.

[Cyrosly]

钢做了个小测试：贴上测试代码，更详细的注意点可以到我的BLOG去看下，代码：

#include<stdio.h>
#include<stdlib.h>
#include<memory.h>
#include<cuda.h>

#pragma comment(lib,"cuda.lib")

int main(int argc,char** argv)
{
	CUdevice dev;
	CUcontext ctx;

	cuInit(0);
	cuDeviceGet(&dev,0);
        cuCtxCreate(&ctx,CU_CTX_SCHED_AUTO,dev);

	CUdeviceptr dpDst,dpSrc;
	cuMemAlloc(&dpDst, 50* 50*sizeof(int));
	cuMemAlloc(&dpSrc,100*100*sizeof(int));

	int* hPtr=new int[10000];

	for(int y=0;y<100;++y){
		for(int x=0;x<100;++x){
			int idx=100*y+x;
			hPtr[idx]=idx;
		}
	}
	cuMemcpyHtoD(dpSrc,hPtr,10000*sizeof(int));
	delete [] hPtr;

	CUDA_MEMCPY2D planeMem;
	memset(&planeMem,0,sizeof(planeMem));
	planeMem.srcMemoryType=CU_MEMORYTYPE_DEVICE;
	planeMem.srcDevice    =dpSrc;
	planeMem.srcXInBytes  =25*sizeof(int);
	planeMem.srcY         =25;
	planeMem.srcPitch     =100*sizeof(int);
	planeMem.dstMemoryType=CU_MEMORYTYPE_DEVICE;
	planeMem.dstDevice    =dpDst;
	planeMem.dstXInBytes  =0;
	planeMem.dstY         =0;
	planeMem.dstPitch     =50*sizeof(int);
	planeMem.WidthInBytes =planeMem.dstPitch;
	planeMem.Height       =50;

	cuMemcpy2DUnaligned(&planeMem);

	FILE* fp=fopen("output.txt","wb");

	hPtr=new int[2500];	
	cuMemcpyDtoH(hPtr,dpDst,2500*sizeof(int));
	for(int y=0;y<50;++y){
		for(int x=0;x<50;++x){
		    fprintf(fp,"%d ",hPtr[50*y+x]);
		}
		fprintf(fp,"\r\n");
	}
    delete [] hPtr;
	fclose(fp);
	
	cuMemFree(dpDst);
	cuMemFree(dpSrc);

	return 0;
}