cuda nbody моделирование - проблема с общей памятью - PullRequest
Новая
       9

cuda nbody моделирование - проблема с общей памятью

1 голос
/ 30 августа 2011

На основе примера вычислительного SDK Nvidia GPU я создал два ядра для симуляции 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;
}

Ответы [ 2 ]

3 голосов
/ 30 августа 2011

Единственный действительно полезный ответ будет получен путем тщательного профилирования, и это только то, что вы можете сделать. NVIDIA поставляет полезные инструменты профилирования для Linux и Windows, сейчас самое время их использовать.

Сказав это, потребление регистров версии с общей памятью значительно больше, чем у версии без общей памяти (37 против 29 при компиляции в цель sm_20 с помощью компилятора выпуска CUDA 4.0). Это может быть простая разница в загруженности, которая вызывает изменение производительности, которое вы видите.

2 голосов
/ 31 августа 2011

На самом деле не разделяемая версия ядра использует разделяемую память в виде кеша L1.Из кода видно, что потоки попадают в одни и те же области глобальной памяти, поэтому они кэшируются и используются повторно.Когда мы добавляем лучшую загруженность и отсутствие дополнительных инструкций (синхронизация и т. Д.), Мы получаем более быстрое ядро.

...