为何加了shared memory后，反而结果精度低，不正确了

加速计算 CUDA
1

我是一步步在矩阵相加相乘的基础上发展的一个小程序。。

原来没有用共享存储器，结果和cpu计算结果相符，但是加入共享存储器之后，在最后几位结果上和CPU的结果有出入

这是美军用共享时的结果
[attach]662124[/attach]

以下是加了共享的结果
[attach]662123[/attach]

如果有需要，可以贴源代码的

2

能贴上源代码吗？

另！这点差别应当是正常的，呵呵！

[ 本帖最后由 yyfn风辰 于 2010-4-30 09:29 编辑 ]

3

这是我的kernel的代码，用了共享存储器的在代码后面加了说明

#ifndef SAMPLE_KERNEL_H_
#define SAMPLE_KERNEL_H_

//**** the following is just for testing ***************
#include<stdio.h>
//******************************************************


/************************************************************************/
/* kernel is beginning*/

__device__ float3 bodyBodyInteraction(float4 bi, float4 bj, float3 ai);
__device__ float3 calculate(float4 a,float4 b,float3 accl,int n);


__device__ float3 bodyBodyInteraction(float4 bi, float4 bj, float3 ai)
{
	float3 r;
	
	// r_ij [3 FLOPS]
	r.x = bj.x - bi.x;
	r.y = bj.y - bi.y;
	r.z = bj.z - bi.z;
	
	// distSqr = dot(r_ij, r_ij) + EPS^2 [6 FLOPS]
	float distSqr = r.x * r.x + r.y * r.y + r.z * r.z + EPS;
	// invDistCube =1/distSqr^(3/2) [4 FLOPS (2 mul, 1 sqrt, 1 inv)]
	
	float distSixth = distSqr * distSqr * distSqr;
	float invDistCube = 1.0f/sqrtf(distSixth);
	
	// s = m_j * invDistCube [1 FLOP]
	float s = bj.w * invDistCube;
	
	// a_i = a_i + s * r_ij [6 FLOPS]
	ai.x += r.x * s;
	ai.y += r.y * s;
	ai.z += r.z * s;
	return ai;
}

__device__ float3 calculate(float4 a,float3 accl,int n)
{
	extern __shared__ float4 Sh_pos[];//added
	int i=0;
	if(i<n)
	{
		accl =bodyBodyInteraction(a,Sh_pos[i],accl);//edited
		i++;
	}
	return accl;
}

__global__ static void GPU(float4* pos,float4* vel,int n)
{
	__shared__ float4 Sh_pos[N];//added
	__shared__ float4 Sh_vel[N];//added

	const int tid=threadIdx.x;
	const int bid=blockIdx.x;
	const int index=bid*blockDim.x+tid;
	float4 atom=pos[index];
	Sh_pos[index]=pos[index];//edited
	Sh_vel[index]=vel[index];//added
	__syncthreads();//added  make sure that all of the data has been loaded

   float3 acc = {0.0f, 0.0f, 0.0f};

	acc=calculate(atom,acc,n);//edited

	/******** added *****************************************************/
	//printf("acc[%d]:%f||%f||%f\n",index,acc.x,acc.y,acc.z);
   
   Sh_vel[index].x +=(float)( acc.x * deltaTime);
   Sh_vel[index].y +=(float)( acc.y * deltaTime);
   Sh_vel[index].z +=(float)( acc.z * deltaTime); 

   Sh_vel[index].x *= damping;
   Sh_vel[index].y *= damping;
   Sh_vel[index].z *= damping;
	//printf("vel[%d]:%f||%f||%f\n",index,vel[index].x,vel[index].y,vel[index].z);
   
   // new position = old position + velocity * deltaTime
   Sh_pos[index].x += Sh_vel[index].x * deltaTime;
   Sh_pos[index].y += Sh_vel[index].y * deltaTime;
   Sh_pos[index].z += Sh_vel[index].z * deltaTime;

	//copy data to global memory
	__syncthreads();
	pos[index].x=Sh_pos[index].x;
	pos[index].y=Sh_pos[index].y;
	pos[index].z=Sh_pos[index].z;
	vel[index].x=Sh_vel[index].x;
	vel[index].y=Sh_vel[index].y;
	vel[index].z=Sh_vel[index].z;
	//printf("pos[%d]:%f||%f||%f\n",index,pos[index].x,pos[index].y,pos[index].z);
/******************************************************************************************************/


}

#endif //  #ifndef SAMPLE_KERNEL_H_
4

呵呵，这样的精度差别可以接受不？？正常不？

5

float distSqr = r.x * r.x + r.y * r.y + r.z * r.z + EPS;
中的EPS是一个很小的数的意思吗？在这个程序中有什么作用呢？

6

能够接受，我写程序的时候也经常出现这个问题，主要是六位小数的后三位不一样，呵呵！

另
Sh_pos[index]=pos[index];//edited
Sh_vel[index]=vel[index];//added

是否有点问题？

7

作用就是防止当r=0的时候后面求倒数出错
呵呵，这个就能过度和我后面判断结果是否对的精度是一样的

8

这两句的意思就是将数据全部拷贝带共享存储器里面

有问题吗？？？？？？？？

9

如果只有一个ｂｌｏｃｋ，应当没有。但是你的这个模式应当是有问题的，ｓｈａｒｅｄ是ｂｌｏｃｋ共享的，呵呵

[ 本帖最后由 yyfn风辰 于 2010-5-2 11:05 编辑 ]

10

呵呵，喔，明白了，谢谢

11

原来如此

12

有不同是正常的，浮点运算是不可结合的
精度是不是异常还是要做一下分析

13

但是我这里是模拟运行啊？？

还没有用到GPU就成现在这样了啊。。。

14

浮点运算是不满足结合率的, 不决定于硬件的微架构