cuda nbody 시뮬레이션 - 공유 메모리 문제

Question

Nvidia GPU 컴퓨팅 SDK의 예제를 기반으로 nbody 시뮬레이션을 위해 두 개의 커널을 만들었습니다. 공유 메모리를 사용하지 않는 첫 번째 커널은 공유 메모리를 사용하는 두 번째 커널보다 ~ 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; 
} 

Answer 1

사실 공유되지 않는 커널 버전은 공유 메모리를 L1 캐시 형태로 사용합니다. 코드에서 스레드는 전역 메모리의 동일한 영역을 차지하므로 캐시되고 재사용됩니다. 우리가 더 많은 점유와 추가 명령 (동기화 등)을 추가 할 때 우리는 더 빠른 커널을 얻습니다.

Answer 2

유일한 정말 유용한 답변이 조심 프로파일에 의해 얻을 수, 그리고는 당신이 할 수있는 위치에있는 무언가이다. NVIDIA는 Linux와 Windows 모두에 유용한 프로파일 링 도구를 제공하며, 이제는이를 사용할 시간이 될 것입니다.

공유 메모리 버전의 레지스터 사용량은 비공유 메모리 버전 (CUDA 4.0 릴리스 컴파일러를 사용하여 sm_20 타겟으로 컴파일 할 때 37 대 29)보다 훨씬 큽니다. 그것은 당신이보고있는 성능의 변화를 일으키는 인원의 단순한 차이 일 수 있습니다.