2

GTX680 でいくつかのパフォーマンス CUDA テストを行っていますが、以下のパフォーマンス結果が得られた理由を理解するのに役立つかどうか疑問に思っていました。私が実行しているコードは次のとおりです。 

#include <stdio.h>
using namespace std;


__global__ void test_hardcoded(int rec,int * output)
{

    int a;
    int rec2=rec/2;
    if(threadIdx.x==1000) *output=rec;
    if(threadIdx.x==1000) *(output+1)=rec2;

    for (int i=0;i<10000;i++)
    {
        __syncthreads();
        a+=i;
    }
    if(threadIdx.x==1000) *output=a;   //will never happen but should fool compiler as to not skip the for loop

}
__global__ void test_softcoded(int rec,int * output)
{
    int a;
    int rec2=rec/2; //This should ensure that we are using the a register not constant memory
    if(threadIdx.x==1000) *output=rec;
    if(threadIdx.x==1000) *(output+1)=rec2;

    for (int i=0;i<=rec2;i++)
    {    __syncthreads();
        a+=i;
    }
    if(threadIdx.x==1000) *output=a;   //will never happen but should fool compiler as to not skip the for loop

}

int main(int argc, char *argv[])
{
    float timestamp;
    cudaEvent_t event_start,event_stop;
    // Initialise
    cudaSetDevice(0);

    cudaEventCreate(&event_start);
    cudaEventCreate(&event_stop);
    cudaEventRecord(event_start, 0);
    dim3 threadsPerBlock;
    dim3 blocks;
    threadsPerBlock.x=32;
    threadsPerBlock.y=32;
    threadsPerBlock.z=1;
    blocks.x=1;
    blocks.y=1000;
    blocks.z=1;

    cudaEventRecord(event_start);
    test_hardcoded<<<blocks,threadsPerBlock,0>>>(10000,NULL);
    cudaEventRecord(event_stop, 0);
    cudaEventSynchronize(event_stop);
    cudaEventElapsedTime(&timestamp, event_start, event_stop);
    printf("test_hardcoded() took  %fms \n", timestamp);

    cudaEventRecord(event_start);
    test_softcoded<<<blocks,threadsPerBlock,0>>>(20000,NULL);
    cudaEventRecord(event_stop, 0);
    cudaEventSynchronize(event_stop);
    cudaEventElapsedTime(&timestamp, event_start, event_stop);
    printf("test_softcoded() took  %fms \n", timestamp);

}

コードによると、私は 2 つのカーネルを実行しています。それらが行うことは、ループして追加することだけです。唯一の違いは、test_softcoded() ループがレジスタと比較するのに対し、test_hardcoded() ループはハードコードされた整数と直接比較することです。

上記のコードを実行すると、次の結果が得られます

$ nvcc -arch=sm_30 test7.cu
$ ./a.out

test_hardcoded() took  51.353985ms 
test_softcoded() took  99.209694ms

test_hardcoded() 関数は、test-softcoded() よりも 2 倍高速です!!!!

test_softcoded() には書き込み後のレジストリ依存関係の可能性があることは理解していますが、レジストリのレイテンシは占有率が高いために完全に隠されているため、非常に高いはずです)、何が問題で、何をすべきか疑問に思っていますtest_softcoded() のパフォーマンスを向上させるために行うこと。

4

1 に答える 1

1

このハードコードされた値により、コンパイラはループ展開などの最適化を行うことができ、パフォーマンスがある程度向上する可能性があります。それが理由かもしれません。

'for (int i=0;i<=rec2;i++)' の前に '#pragma unroll 5000' のようなコードを追加して実行すると、疑問が解決します。

于 2012-10-22T16:25:58.880 に答える