2

幅と高さが使用しているブロックサイズの倍数である 2 つの行列を乗算する CUDA カーネルがあります。

私が使用している Nvidia Quadro Fx 3800 の理論上の帯域幅は 50 Gb/s で、奇妙な結果が得られています (実効帯域幅が理論上の帯域幅よりも大きい)。

ここにいくつかの結果を投稿します。

ブロックサイズ 2 あり

[10][10] * [10][10] -> 帯域幅=0,02 Gb/秒 [1000][1000]*[1000][1000] -> 帯域幅=69,4 Gb/秒

ブロックサイズ 64 の場合

[1000][1000] * [1000][1000] -> 帯域幅=486,4 Gb/秒 [10000][10000] * [10000][10000] -> 帯域幅= 45072,12 Gb/秒

Nvidia のベスト プラクティス ガイドから有効な帯域幅の式を使用しました (簡略化していますが、同等のものです (愚かな間違いがない限り))。カーネルは、私が読んだいくつかのNvidia Lecturesと非常に似ているため（等しくない場合）、また適切に機能しているため（afaik）、問題ないと思います。

#define blocksize 64
#define HM (10000) 
#define WM (10000) 
#define WN (10000)
#define HN WM 
#define WP WN   
#define HP HM  
#define PTH WM
#define PTW HM

__global__ void nonsquare(float*M, float*N, float*P, int uWM,int uWN)
   {
__shared__ float MS[blocksize][blocksize];
__shared__ float NS[blocksize][blocksize];

int tx=threadIdx.x, ty=threadIdx.y, bx=blockIdx.x, by=blockIdx.y;
int rowM=ty+by*blocksize;
int colN=tx+bx*blocksize;
int Pvalue=0;

for(int m=0; m< uWM/blocksize;m++){
    MS[ty][tx]=M[rowM*uWM+(m*blocksize+tx)];
    NS[ty][tx]=M[colN + uWN*(m*blocksize+ty)];
    __syncthreads();
    for(int k=0;k<blocksize;k++)
        Pvalue+=MS[ty][k]*NS[k][tx];
    P[rowM*WP+colN]=Pvalue;
}

}
int main(){


cudaEvent_t evstart, evstop;
cudaEventCreate(&evstart);
cudaEventCreate(&evstop);

float*M=(float*)malloc(sizeof(float)*HM*WM);
float*N=(float*)malloc(sizeof(float)*HN*WN);

for(int i=0;i<WM*HM;i++)
    M[i]=(float)i;
for(int i=0;i<WN*HN;i++)
    N[i]=(float)i;




float*P=(float*)malloc(sizeof(float)*HP*WP);



float *Md,*Nd,*Pd;
cudaMalloc((void**)&Md,HM*WM*sizeof(float));

cudaMalloc((void**)&Nd,HN*WN*sizeof(float));

cudaMalloc((void**)&Pd,HP*WP*sizeof(float));



cudaMemcpy(Md,M,HM*WM*sizeof(float),cudaMemcpyHostToDevice);

cudaMemcpy(Nd,N,HN*WN*sizeof(float),cudaMemcpyHostToDevice);



dim3 dimBlock(blocksize,blocksize);//(tile_width , tile_width);
dim3 dimGrid(WN/dimBlock.x,HM/dimBlock.y);//(width/tile_width , width/tile_witdh);

cudaEventRecord(evstart,0);

nonsquare<<<dimGrid,dimBlock>>>(Md,Nd,Pd,WM,WN);

cudaEventRecord(evstop,0);
cudaEventSynchronize(evstop);
float time;
cudaEventElapsedTime(&time,evstart,evstop);

cudaMemcpy(P,Pd,WP*HP*sizeof(float),cudaMemcpyDeviceToHost);

    cudaFree(Md);
cudaFree(Nd);
cudaFree(Pd);


    printf("\ntime spent:%f",time);
float Bandwidth=(HM*WM*4+WN*HN*4+HP*WP*4)/(time*1000000); /
printf("\nEffective Bandwidth:%f Gb/s\n",Bandwidth);
    }

前もって感謝します

4

3 に答える 3

2

カーネルが静かに失敗しているだけだと思います。

  1. カーネルの呼び出し後にエラーがないか確認しましたか?

  2. コードは機能しますか?

  3. タイミングでどのような結果がありますか?

于 2011-03-28T06:17:33.863 に答える
1

共有メモリやテクスチャメモリなどを使用すると、理論上の帯域幅を超える可能性があることに注意してください。これは、ハードウェアでサポートされている専用の機能（組み込みの双一次テクスチャ補間など）を、おそらく意図せずに利用していることを意味します。

Robert Harveyが言及した理由に加えて、ベンダーによるカードの工場でのオーバークロックも発生する可能性があります（ただし、QuadrosよりもGeForceの方が一般的です）。

全体として、（メモリまたはコンピューティングのいずれかで）理論上の帯域幅に近づいたり超えたりした場合は、うまくいっていると思います。

于 2011-03-23T13:02:43.620 に答える
0

私はいくつかの説明を考えることができます：

  1. 測定に悪影響を与えるベースライン コードの変更
  2. 無効なパフォーマンスの仮定
  3. 未確認のマイクロ最適化。
  4. 非現実的なベンチマーク。

あなたのコードは単純化されていると言います。元のベンチマーク コードを使用してみて、何が起こるか見てみます。数値がより現実的である場合は、元のベンチマーク コードと単純化されたコードを比較して違いを特定できます。

于 2011-03-22T19:08:23.210 に答える