ピークパフォーマンスに達することができません

Question

それぞれのピークパフォーマンスに到達しようとしています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）

Answer 1

私は正しいFLOP計算を用いて、一緒にあなたのコードから完全にREPROケースを置く場合：

#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で実行してください：

純粋なFMADコード（1.4 GHz * 2 FLOP * 8コア/ MP * 30 MP）= 672 GFLOP/sを実行しているカードのピークFLOP/sの0.01％以内になります。

実際には、コードはマルチプロセッサあたり1ブロックでピークFLOP/sに達しますが、FLOP/s数は正しく計算されていないようです。