Не удается достичь максимальной производительности - PullRequest
Новая
       10

Не удается достичь максимальной производительности

2 голосов
/ 03 января 2012

Я пытаюсь достичь максимальной производительности каждого SM из приведенного ниже кода.Пик лежит где-то между 25 GFlops (GTX275-GT200 Arch.).Этот код дает максимум 8 GFlops 

__global__ void new_ker(float *x)
{
  int index = threadIdx.x+blockIdx.x*blockDim.x;
  float a,b;
  a=0;
  b=x[index];
  //LOOP=10000000
  //No. of blocks = 1
  //Threads per block = 512 (I'm using GTX 275 - GT200 Arch.)
  #pragma unroll 2048
  for(int i=0;i<LOOP;i++){
       a=a*b+b;
  }  

  x[index] = a;

 }

Я не хочу увеличивать ILP в коде.Есть идеи, почему он не достигает пика ?? 

int main(int argc,char **argv)
{

   //Initializations
   float *x;
   float *dx;
   cudaEvent_t new_start,new_stop;
   float elapsed;
   double gflops;
   x = 0;
   flag = 0;
   cudaMalloc((void **)&dx,sizeof(float)*THPB);

   //ILP=1  
   cudaEventCreate(&new_start);
   cudaEventCreate(&new_stop);
   printf("Kernel1:\n");
   cudaEventRecord(new_start, 0);
   new_ker<<<BLOCKS,THPB>>>(dx);
   cudaEventRecord(new_stop,0);
   cudaEventSynchronize(new_stop);
   cudaEventElapsedTime(&elapsed,new_start,new_stop);
   x = (float *)malloc(sizeof(float)*THPB);
   cudaMemcpy(x,dx,sizeof(float)*THPB,cudaMemcpyDeviceToHost);

   gflops = ((double)(BLOCKS)*(THPB)*LOOP/elapsed)/1000000;
   printf("\t%f",gflops);
   cudaEventDestroy(new_start);
   cudaEventDestroy(new_stop);
   return 0;
}

Платформа: CUDA 3.0 NVIDIA GeForce GTX275 (GT200)

1 Ответ

4 голосов
/ 03 января 2012

Если я соберу полный сценарий воспроизведения из вашего кода, используя правильный расчет FLOP: 

#include <stdio.h> 

#define LOOP (10000000)
#define BLOCKS (30)
#define THPB (512)

__global__ void new_ker(float *x)
{
  int index = threadIdx.x+blockIdx.x*blockDim.x;
  float a,b;
  a=0;
  b=x[index];
  #pragma unroll 2048
  for(int i=0;i<LOOP;i++){
       a=a*b+b;
  }  

  x[index] = a;
}

int main(int argc,char **argv)
{

   //Initializations
   float *x;
   float *dx;
   cudaEvent_t new_start,new_stop;
   float elapsed;
   double gflops;
   x = 0;
   cudaMalloc((void **)&dx,sizeof(float)*THPB);

   //ILP=1  
   cudaEventCreate(&new_start);
   cudaEventCreate(&new_stop);
   printf("Kernel1:\n");
   cudaEventRecord(new_start, 0);
   new_ker<<<BLOCKS,THPB>>>(dx);
   cudaEventRecord(new_stop,0);
   cudaEventSynchronize(new_stop);
   cudaEventElapsedTime(&elapsed,new_start,new_stop);
   x = (float *)malloc(sizeof(float)*THPB*BLOCKS);
   cudaMemcpy(x,dx,sizeof(float)*THPB*BLOCKS,cudaMemcpyDeviceToHost);

   gflops = 2.0e-6 * ((double)(LOOP)*double(THPB*BLOCKS)/(double)elapsed);
   printf("\t%f\n",gflops);
   cudaEventDestroy(new_start);
   cudaEventDestroy(new_stop);
   return 0;
}

И я скомпилирую его и запусту на 1,4 ГГц GTX275 с CUDA 3.2 на 64-битной Linuxплатформа: 

$ nvcc -arch=sm_13 -Xptxas="-v" -o perf perf.cu
ptxas info    : Compiling entry function '_Z7new_kerPf' for 'sm_13'
ptxas info    : Used 4 registers, 8+16 bytes smem, 8 bytes cmem[1]
$ ./perf 
Kernel1:
        671.806039

Я получаю в пределах 0,01% от максимальной FLOP / с для этой карты, использующей чистый код FMAD (1,4 ГГц * 2 FLOP * 8 ядер / MP * 30 MP) = 672 GFLOP / с.

Таким образом, кажется, что код фактически достигает пика FLOP / s с одним блоком на мультипроцессор, но вы просто неправильно вычисляете число FLOP / s.

...