无法达到最佳性能

问题描述

我正在尝试通过下面的代码达到每个 SM 的最高性能.最高峰位于25 GFlops(GTX275-GT200 Arch.)之间.该代码最多可提供8个GFlops.

I'm trying to reach peak performance of each SM from the code below. The peak lies somewhere between 25 GFlops(GTX275-GT200 Arch.). This code gives 8 GFlops at the max.

__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.为什么它没有达到顶峰的任何想法?

I don't want to increase ILP in the code. Any ideas why it's not reaching peak??

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.0NVIDIA GeForce GTX275(GT200)

Platform:CUDA 3.0NVIDIA GeForce GTX275 (GT200)

推荐答案

如果我使用正确的FLOP计算方法从您的代码中整理出完整的repro案例:

If I put together a complete repro case from your code, using the correct FLOP calculation:

#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;
}

我将其编译并在64位linux平台上的带有CUDA 3.2的1.4GHz GTX275上运行:

And I compile it and run it on a 1.4GHz GTX275 with CUDA 3.2 on a 64 bit linux platform:

$ 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

对于运行纯FMAD代码(1.4 GHz * 2 FLOP * 8核心/MP * 30 MP)= 672 GFLOP/s的那张卡，我得到的峰值FLOP/s不到0.01％.

I get within 0.01% of peak FLOP/s for that card running a pure FMAD code (1.4 GHz * 2 FLOP * 8 cores/MP * 30 MP) = 672 GFLOP/s.

因此，似乎该代码确实确实达到了每个多处理器一个块的FLOP/s峰值，但是您只是没有正确计算FLOP/s数.

So it seems that the code does, in fact, hit peak FLOP/s with one block per multiprocessor, but you just are not calculating the FLOP/s number correctly.

这篇关于无法达到最佳性能的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持！