0

私はcudaについて勉強してきました。

CUDAプログラミングガイドでは、共有メモリアクセス時間はグローバルメモリ時間よりも高速です。

そこで、行列の乗算を実行するコードを作成しました。

これが私のコードです。version1はグローバルメモリを使用し、version2は共有メモリを使用しました

私のGPUはteslac2070cudasdkバージョン4.2です

メインコード

#define Matrix_Width   9216
#define Matrix_Divide  4
#define Tile_Width     32
#define Ce_Size 4096

#if Matrix_Width == 9216
    #define Matrix_Size 9216*9216
#elif Matrix_Width == 12800
    #define Matrix_Size 12800*12800
#elif Matrix_Width == 15872
    #define Matrix_Size 15872*15872
#elif Matrix_Width == 18432
    #define Matrix_Size 18432*18432
#endif

float* H_Input1 = (float*)malloc( sizeof(float) * Matrix_Size );
float* H_Input2 = (float*)malloc( sizeof(float) * Matrix_Size );
float* H_Output = (float*)malloc( sizeof(float) * Matrix_Size );

for( int i=0 ; i < Matrix_Size ; i++ ){
    H_Input1[i] = 1.0f;
H_Input2[i] = 1.0f;
}
memset( H_Output, 0 , sizeof(float) * Matrix_Size );

float* D_Input1;
float* D_Input2;
float* D_Output;

cudaMalloc( (void**)&D_Input1, sizeof(float) * Matrix_Size );
cudaMalloc( (void**)&D_Input2, sizeof(float) * Matrix_Size );
cudaMalloc( (void**)&D_Output, sizeof(float) * Matrix_Size );
cudaMemcpy( D_Input1, H_Input1, sizeof(float) * Matrix_Size, cudaMemcpyHostToDevice );
cudaMemcpy( D_Input2, H_Input2, sizeof(float) * Matrix_Size, cudaMemcpyHostToDevice );
cudaMemcpy( D_Output, H_Output, sizeof(float) * Matrix_Size, cudaMemcpyHostToDevice );


event_pair Event;
start_timer( &Event );
dim3  dimGrid( Matrix_Width/Matrix_Divide/Tile_Width, Matrix_Width/Matrix_Divide/Tile_Width, 1 );
dim3 dimBlock( Tile_Width, Tile_Width, 1 );

kernel_global<< dimGrid, dimBlock>>>( D_Input1, D_Input2, D_Output );
stop_timer( &Event, "1GB mMemory Test\n" );
cudaMemcpy( H_Output, D_Output, sizeof(float) * Matrix_Size, cudaMemcpyDeviceToHost );

カーネルバージョン1

__global__ void kernel_global( float* Input1, float* Input2, float* Output ){

for( int i = 0 ; i < Matrix_Divide ; i++ ){
for( int j = 0 ; j < Matrix_Divide ; j++ ){

float Sum = 0;
int Row = (i * (Matrix_Width/Matrix_Divide)) + (blockIdx.y * Tile_Width) + threadIdx.y; 
int Col = (j * (Matrix_Width/Matrix_Divide)) + (blockIdx.x * Tile_Width) + threadIdx.x;

    for( int k = 0 ; k < Matrix_Width ; k++ ){
        Sum += Input1[ Row * Matrix_Width + k ] * Input2[ k * Matrix_Width + Col ];
    }
    Output[ Row*Matrix_Width+Col] = Sum;
    }
    }
}

カーネルバージョン2

    __global__ void kernel_shared( float* Input1, float* Input2, float* Output ){

    __shared__ float Input1_s[Tile_Width][Tile_Width];
    __shared__ float Input2_s[Tile_Width][Tile_Width];

    int Bx = blockIdx.x;
    int By = blockIdx.y;
    int Tx = threadIdx.x;
    int Ty = threadIdx.y;

    for( int i = 0 ; i < Matrix_Divide ; i++ ){
        for( int j = 0 ; j < Matrix_Divide ; j++ ){

            float Sum = 0;
            int Row = (i * (Matrix_Width/Matrix_Divide)) + (By * Tile_Width) + Ty;
            int Col = (j * (Matrix_Width/Matrix_Divide)) + (Bx * Tile_Width) + Tx;


            for( int m = 0 ; m < Matrix_Width/Tile_Width ; m++ ){

                Input1_s[Ty][Tx] = Input1[ Row * Matrix_Width + ( m * Tile_Width + Tx ) ];
                Input2_s[Ty][Tx] = Input2[ ( m * Tile_Width + Ty ) * Matrix_Width + Col ];
                __syncthreads();


                for( int k = 0 ; k < Tile_Width; k++ ){
                    Sum += Input1_s[Ty][k] * Input2_s[k][Tx];
                }
                __syncthreads();
            }
            Output[ Row*Matrix_Width+Col] = Sum;
        }
    }
}

このコードは、幅=9216のマトリックスを作成しました

一度に計算することはできません。ブロックの最大数が65535で、スレッドが1024であるため

したがって、私は4を使用して行列の幅を分割したので、行列は16のチャンクを分割します。

一度に1つのチャンクを計算できます。

だから私はloopcountが16であるループを使用しました（i * j = 16）

チャンクはブロックとスレッドに分割されます。（tile_width = 32）

テスト結果はとても奇妙です。

バージョン1は90秒かかりました

バージョン2は130秒かかりました

この結果がわかりません

共有メモリ要素はタイルで再利用されていると思います...

なぜバージョン1がバージョン2よりも速いのですか？

よろしくお願いします！！

4

1 に答える 1

3

Tesla C2070は、グローバルメモリアクセスをキャッシュするコンピューティング機能2.0デバイスです。したがって、どちらの場合も、内部ループは（最初の反復を除いて）オンチップメモリ​​で動作します。

唯一の違いは、共有メモリにデータをロードするための2番目のバージョンの追加コードには余分な時間がかかるのに対し、最初のバージョンでは、この作業はキャッシュロジックによってハードウェアで実行されることです。

共有メモリにデータを手動でキャッシュすることは、コンピューティング機能1.xデバイスでのみ、またはハードウェアキャッシュのLRU（最も最近使用されていない）ロジックよりもどのデータが再利用されているかをよく理解している場合にのみ価値があります。

于 2012-09-21T09:42:42.640 に答える