Cuda无人模拟-共享内存问题
cuda nbody simulation - shared memory problem
基于Nvidia GPU计算SDK的示例,我创建了两个内核用于无体仿真。第一个不使用共享内存的内核比第二个使用共享内存的内核快15%。为什么使用共享内存的内核会变慢?
内核的目的:- 8192具尸体,
- 每块线程数= 128,
- 每个网格块= 64。
- 设备:GeForce GTX 560 Ti.
第一内核:
#define N 8192
#define EPS2 0.001f
__device__ float4 vel[N];
__device__ float3 force(float4 bi, float4 bj, float3 ai)
{
float3 r;
r.x = bj.x - bi.x;
r.y = bj.y - bi.y;
r.z = bj.z - bi.z;
float distSqr = r.x * r.x + r.y * r.y + r.z * r.z + EPS2;
float distSixth = distSqr * distSqr * distSqr;
float invDistCube = 1.0f/sqrtf(distSixth);
float s = bj.w * invDistCube;
ai.x += r.x * s;
ai.y += r.y * s;
ai.z += r.z * s;
return ai;
}
__global__ void points(float4 *pos, float dt)
{
int k = blockIdx.x * blockDim.x + threadIdx.x;
if(k >= N) return;
float4 bi, bj, v;
float3 ai;
v = vel[k];
bi = pos[k];
ai = make_float3(0,0,0);
for(int i = 0; i < N; i++)
{
bj = pos[i];
ai = force(bi, bj, ai);
}
v.x += ai.x * dt;
v.y += ai.y * dt;
v.z += ai.z * dt;
bi.x += v.x * dt;
bi.y += v.y * dt;
bi.z += v.z * dt;
pos[k]=bi;
vel[k]=v;
}
第二个内核:
#define N 8192
#define EPS2 0.001f
#define THREADS_PER_BLOCK 128
__device__ float4 vel[N];
__shared__ float4 shPosition[THREADS_PER_BLOCK];
__device__ float3 force(float4 bi, float4 bj, float3 ai)
{
float3 r;
r.x = bj.x - bi.x;
r.y = bj.y - bi.y;
r.z = bj.z - bi.z;
float distSqr = r.x * r.x + r.y * r.y + r.z * r.z + EPS2;
float distSixth = distSqr * distSqr * distSqr;
float invDistCube = 1.0f/sqrtf(distSixth);
float s = bj.w * invDistCube;
ai.x += r.x * s;
ai.y += r.y * s;
ai.z += r.z * s;
return ai;
}
__device__ float3 accumulate_tile(float4 myPosition, float3 accel)
{
int i;
for (i = 0; i < THREADS_PER_BLOCK; i++)
{
accel = force(myPosition, shPosition[i], accel);
}
return accel;
}
__global__ void points(float4 *pos, float dt)
{
int k = blockIdx.x * blockDim.x + threadIdx.x;
if(k >= N) return;
float4 bi, v;
float3 ai;
v = vel[k];
bi = pos[k];
ai = make_float3(0.0f, 0.0f, 0.0f);
int i,tile;
for(tile=0; tile < N / THREADS_PER_BLOCK; tile++)
{
i = tile * blockDim.x + threadIdx.x;
shPosition[threadIdx.x] = pos[i];
__syncthreads();
ai = accumulate_tile(bi, ai);
__syncthreads();
}
v.x += ai.x * dt;
v.y += ai.y * dt;
v.z += ai.z * dt;
bi.x += v.x * dt;
bi.y += v.y * dt;
bi.z += v.z * dt;
pos[k]=bi;
vel[k]=v;
}
唯一真正有用的答案将通过仔细的分析获得,这是您唯一能够做的事情。NVIDIA为Linux和Windows提供了有用的分析工具,现在可能是时候使用它们了。
话虽如此,共享内存版本的寄存器消耗比非共享内存版本大得多(使用CUDA 4.0版本编译器编译到sm_20目标时为37 vs 29)。这可能是占用率的简单差异导致了您所看到的性能变化。
实际上非共享版本的内核确实以L1缓存的形式使用共享内存。从代码中我们可以看到,线程访问了相同的全局内存区域,因此它被缓存和重用。
相关文章:
- C++ 中 std::vector 的内存问题
- 实现快速排序的内存问题
- 此add_node函数会导致内存问题吗?
- C++:为什么这段代码给我内存问题/未定义的行为?
- 我不明白附加字符串的内存问题
- 是否有可能存在不会崩溃程序的内存问题
- 在发布版本中删除类指针会导致内存问题
- Android OpenCV 应用程序的内存问题
- 线程的内存问题(微小线程,C++)
- 递归结构的向量有内存问题
- C++合并排序内存问题
- 将智能指针发送到Protobaf.内存问题
- 潜在的动态内存问题
- JNI 中的内存问题
- C++Lambda函数关闭-内存问题
- 解决由全局静态变量引起的内存问题
- 类堆栈内存问题(致命错误)
- QNetworkAccessManager内存问题
- 读取进程内存问题未更新
- 类指针内存问题