cuda nbody模擬 - 共享內存問題

基於Nvidia GPU計算SDK的示例，我爲nbody模擬創建了兩個內核。沒有利用共享內存的第一個內核比使用共享內存的第二個內核快了15％。爲什麼共享內存的內核速度較慢？cuda nbody模擬 - 共享內存問題

內核參數：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; 
}

來源

2011-08-30 Lynx Lynx

實際上非共享版本的內核確實使用L1緩存形式的共享內存。從代碼中我們可以看到線程碰到了全局內存的相同區域，所以它被緩存和重用。當我們添加更好的佔用率和缺乏額外的指令（同步等）時，我們會得到更快的內核。

來源

2011-08-31 17:19:42 llukas

唯一真正有用的答案將通過仔細分析可以得到，這是唯一的，你是在一個位置，做一些事情。 NVIDIA推出適用於Linux和Windows的實用分析工具，現在可能是使用它們的時候了。

話雖如此，共享內存版本的寄存器消耗遠遠大於非共享內存版本（使用CUDA 4.0版本編譯器編譯爲sm_20目標時，版本爲37和29）。這可能是一個簡單的差異，導致你看到的表現發生變化。

來源

2011-08-30 19:56:40 talonmies

cuda nbody模擬 - 共享內存問題

回答

相關問題