我又来请求帮忙了．．．

加速计算 CUDA
1 
__global__ void MyKernel(float *arrSd,float *arrEd,float *arrRd,int n)
{
	int i=threadIdx.x;
	int j=threadIdx.y;
	int bid=blockIdx.x*blockDim.x;
	int id=j*blockDim.x+i;
	__shared__ float mid[64][2];
	mid[i][j]=0;
	if(id<n)
	{
		mid[i][j]=arrSd[id]-arrEd[id];
		mid[i][j]=mid[i][j]*mid[i][j];
		
		__syncthreads();

		if(j==0)
		{
			mid[i][j]=mid[i][j]+mid[i][j+1];
			arrRd[i+bid]=sqrtf(mid[i][j]);
		}
	}
}


int main()
{
	cudaEvent_t sd,ed;


	int streamSize=512;
	const int n=1024;	

	float *arrS=0,*arrE=0,*arrR=0;	
	arrS=(float *)malloc(sizeof(float)*n);
	arrE=(float *)malloc(sizeof(float)*n);
	arrR=(float *)malloc(sizeof(float)*n);

	const int nstreams=2;	
	int size=sizeof(float)*n;

	for(int i=0;i!=n;i++)
	{
		arrS[i]=i+1;
		arrE[i]=i;
		arrR[i]=0;
	}

	cudaStream_t stream[nstreams];
	for(int i=0;i!=nstreams;++i)
		cudaStreamCreate(&stream[i]);

	float *arrSd,*arrEd,*arrRd;


	cudaMalloc(&arrSd,size);
	cudaMalloc(&arrEd,size);
	cudaMalloc(&arrRd,size/2);

	dim3 threads,blocks;
	
	threads=dim3(64,2);

	blocks=dim3(streamSize/(2*threads.x),1);
	int sharesize=sizeof(float)*threads.x*2;
	

	cudaEventCreate(&sd);
	cudaEventCreate(&ed);
	cudaEventRecord(sd,0);

	for(int i=0;i!=nstreams;++i)
	{
		cudaMemcpyAsync(arrSd+i*streamSize,arrS+i*streamSize,sizeof(float)*streamSize,cudaMemcpyHostToDevice,stream[i]);

		cudaMemcpyAsync(arrEd+i*streamSize,arrE+i*streamSize,sizeof(float)*streamSize,cudaMemcpyHostToDevice,stream[i]);

		MyKernel<<<blocks,threads,sharesize,stream[i]>>>(arrSd+i*n/nstreams,arrEd+i*n/nstreams,arrRd+i*n/(nstreams*2),n);

		cudaMemcpyAsync(arrR+i*streamSize/2,arrRd+i*streamSize/2,sizeof(float)*streamSize/2,cudaMemcpyDeviceToHost,stream[i]);
	}
	cudaDeviceSynchronize();
	cudaEventRecord(ed,0);
	cudaEventSynchronize(ed);

	float tid;
	cudaEventElapsedTime(&tid,sd,ed);
	for(int i=0;i!=nstreams;i++)
		cudaStreamDestroy(stream[i]);
	

	for(int i=0;i!=n/2;i++)
	{
		printf("%f ",arrR[i]);
	}
	
	printf(",%f,%f",ti,tid);

	cudaEventDestroy(sd);
	cudaEventDestroy(ed);
	cudaFree(arrSd);
	cudaFree(arrEd);
	cudaFree(arrRd);

	free(arrS);
	free(arrE);
	free(arrR);

	int i=0;
	scanf("%d",&i);
}

其实，我这个是毕业设计，要求实现向量距离的大规模并行计算，现在做成这样，但是加上数据传输怎么都比CPU慢…我已经不知道怎么办了…请各位达人支支招.

还有想问那个zero copy到底是怎么用的，我一直看不懂，可以和流一起用吗？除此之外还有别的办法么

2

LZ您好：

大致看了您的代码，简要说明如下：

1：您的计算规模太小，这个规模比CPU慢十分正常。

2：您的stream使用无法实现计算和传输相掩盖，因为您使用了malloc申请的主机端的线性缓冲区，而不是使用cudaHostAlloc()申请的page-locked memory。以及，请尽量使用telsa系列的显卡以保证计算和掩盖能顺利进行。

3：您的kernel使用shared memory有bank conflict。

4：以及您的kernel可能有更优化的写法。

5：您使用了kernel内静态分配的shared memory，这样调用kernel的时候就无需指定动态划分shared memory的空间大小。指定了会浪费空间，但只要shared memory够用，无其他副作用。

6：关于zero-copy，您可以试用一下，毕竟您的数据只需要读入一次，回写一次。您可以参考programming guide和其他讨论帖的内容。以及可以和流一起用，以及未必需要和流一起用。以及您也可以先只考虑多stream掩盖传输时间。

7：实际上要和CPU相比的话，您需要评估一下您的数据量，计算速度和数据传输速度。如果单纯传输已经超过了CPU用时，那么改写为GPU代码是无意义的。以及如果数据量太小，也是没有什么意义的，请用CPU直接搞定。

大致回答如上，祝您编码顺利~

3

谢谢！ 我会根据你的建议试试的！

4

不客气的，欢迎您继续尝试并回馈论坛~

祝您好运~

5

楼主您好，我将从问题规模，计算量的大小，传输的代价，以及对比现在的CPU的多个方面为您进行“理论上的”分析：

楼主您计算的是: sqrt((An - An+1)^2 + (Bn - Bn+1)^2))。对于特定的n属于[0,1024]，您每个n需要3次加减法，2次乘法，1次sqrt, 一共6次运算。

而每次您平均是4个float进去，1个float结果出来，平均是5个float的读写对比6次计算。如果sqrtf是 SFU算的，将需要2条指令，基本上卡在5float读写 : 2 SFU指令上。
在3.0这种卡上，（1组LSU, 1组SFU), 将卡在读写显存上。此时您的kernel的性能将只受限于您的显存的读写带宽。

然后因为您的数据基本上没有重用，基本上复制到显存(PCI-E受限)，显存读写(显存带宽受限), 就复制回来了(PCI-E受限),
对于您的N组数据，您将最短需要8N字节/PCI-E带宽 + 12N字节/显存带宽 + 4N字节 / PCI-E带宽， 也就是您将至少需要：

12N字节 / PCI-E带宽 + 12N字节/显存带宽.

而您如果在CPU上算这个，对于现在的CPU, 6次加法、减法、乘法、平方根，都可以用SSE计算，而且稍微组合， 一半的ivy bridge之类的CPU, 都能取得24次float计算/cycle, 此时CPU也将卡在访问内存上。此时的CPU需要的时间将是：
12N字节/(内存带宽)

因为现在往往内存带宽一半在20GB/S左右(DDR3-1600, 双通道), 远比上文的时间的前一半的PCI-E带宽大，您直接在CPU上要快的多。

此外，因为还要有传输开销，和显卡交互（例如启动kernel)的开销，在您这么小的数据规模（12KB数据）下，这些代销本身甚至都可能比您的kernel运行要长，您的kernel的平均时间再慢了基本无问题。

所以您在CPU上，可能会快几倍的。

建议此问题, 在CPU上算。
不要为了CUDA而CUDA, CPU上计算这个，也不过只是一次读写的时间。用CUDA无意义了。

6

给楼主个直观的数据，

您12KB的数据量，一半的很烂的几百元的APU, 单通道DDR3(~10GB/S访存), 只需要1us!!

而您想想您用CUDA, 从头到尾，1us您能干嘛？您1us估计数据都没传到device上呢。

是吧。

7

很谢谢版主的分析，但是很可惜，因为我的毕业设计．．就是要用CUDA，并且要它与CPU对比，这已经是不能改变的事实了．．．所以我没办法．．．不过还是谢谢楼主了，每次都很认真的帮我们解决问题．

8

既然楼主一定要用，那么我也没办法。

上文说的很清楚了，您直接就地计算的时间，也远小于您传输的时间。

我没有更好的建议了。

9

LZ您好，您或许可以更换一个用GPU完成的其他更适合GPU的题目，这样不违背您的需求，也能体现出合适的GPU加速结果。

祝您好运~

10

诶…这也不是我可以改变的…最初我是不懂CUDA，完了老师说很简单的，你只要学点东西就好了…然后就…这些题目和要求都是老师固定了的…

11

你可以加大数据量，将算法换成复杂点的，就会好很多。

不妨试试改改看，既然你BOSS糊弄你，你就糊弄他吧。

原来的直接了当的写，还是建议CPU. 但您可以糊弄的，不是吗？

12

那你也可以反驳他，并做另外一个超过他水平的东西出来，并以此震慑贵BOSS。

祝您马到成功！

13

…版主大人们…你两太凶猛了

14

…版主大人们…你两太凶猛了