-3 
// Includes
#include <stdio.h>
#include <cutil_inline.h>
#include <shrQATest.h>
#include <time.h>
#define CLOCKS_PER_SEC ((clock_t)1000)

// Variables
float* h_A;
float* h_B;
float* h_C;
float* h_C_cpu;
float* d_A;
float* d_B;
float* d_C;
bool noprompt = false;

// Functions
void CleanupResources(void);
void RandomInit(float*, int);
void ParseArguments(int, char**);
void ZeroInit(float*, int);
// Device code

__global__ void MatrixMul(const float*A,const float*B,float*C,int Arow,int Acol,int Bcol)
{
    int coli= blockDim.x * blockIdx.x + threadIdx.x;
    int rowi= blockDim.y * blockIdx.y + threadIdx.y;
    float tmp=0;
    C[rowi*Bcol+coli]=0;
    for(int k=0;k<Acol;k++)
    {
        if(rowi<Arow&&coli<Bcol)
            C[rowi*Bcol+coli]+=A[rowi*Acol+k]*B[k*Bcol+coli];
    }
    //__syncthreads();
    //C[rowi*Bcol+coli]=tmp;
}

// Host code
int main(int argc, char** argv)
{
    shrQAStart(argc, argv);
    clock_t start,end;
    double duration;
    printf("Vector Addition\n");
    int a_row=800,a_col=600,b_row=600,b_col=900;
    int a_size =a_row*a_col* sizeof(float);
    int b_size=b_row*b_col*sizeof(float);
    int c_size=a_row*b_col*sizeof(float);
    //const int matrixrow=10000,matrixcol=10000;


    h_A=(float*)malloc(a_size);
    h_B=(float*)malloc(b_size);
    h_C=(float*)malloc(c_size);
    h_C_cpu=(float*)malloc(c_size);
    RandomInit(h_A, a_size/sizeof(float));
    RandomInit(h_B, b_size/sizeof(float));
    //memset(h_C,0,c_size);
    ZeroInit(h_C,c_size/sizeof(float));
    //memset(h_C_cpu,0,c_size);
    ZeroInit(h_C_cpu,c_size/sizeof(float));
    //RandomInit(h_C, c_size);
    start=clock();
    int i,j,k;
    for(i=0;i<a_row;i++)
    {
        for(j=0;j<b_col;j++)
        {
            for(k=0;k<a_col;k++)
            {
                h_C_cpu[i*b_col+j]+=h_A[i*a_col+k]*h_B[k*b_col+j];
            }
        }
    }
    end=clock();
    duration=double(end-start)/CLOCKS_PER_SEC;
    printf("CPU time: %lf\n",duration);



    cutilSafeCall(cudaMalloc((void**)&d_A,a_size));
    cutilSafeCall(cudaMalloc((void**)&d_B,b_size));
    cutilSafeCall(cudaMalloc((void**)&d_C,c_size));


    ParseArguments(argc, argv);

    // Allocate input vectors h_A and h_B in host memory
    /*h_A = (float*)malloc(size);
    if (h_A == 0) CleanupResources();
    h_B = (float*)malloc(size);
    if (h_B == 0) CleanupResources();
    h_C = (float*)malloc(size);
    if (h_C == 0) CleanupResources();*/

    // Initialize input vectors


    // Allocate vectors in device memory
    /*cutilSafeCall( cudaMalloc((void**)&d_A, size) );
    cutilSafeCall( cudaMalloc((void**)&d_B, size) );
    cutilSafeCall( cudaMalloc((void**)&d_C, size) );*/
    start=clock();
    // Copy vectors from host memory to device memory
    cutilSafeCall( cudaMemcpy(d_A, h_A, a_size, cudaMemcpyHostToDevice) );
    cutilSafeCall( cudaMemcpy(d_B, h_B, b_size, cudaMemcpyHostToDevice) );

    // Invoke kernel
    //int threadsPerBlock = 1024;
    dim3 dimblock(32,32);
    int blockx = (b_col + dimblock.x - 1) /dimblock.x;
    int blocky = (a_row + dimblock.y - 1) /dimblock.y;
    dim3 dimgrid(blockx,blocky);

    MatrixMul<<<dimgrid, dimblock>>>(d_A,d_B,d_C,a_row,a_col,b_col);


    //myVecAdd<<<1,threadsPerBlock>>>(d_A,d_B,d_C,N);
    cutilCheckMsg("kernel launch failure");
#ifdef _DEBUG
    cutilSafeCall( cutilDeviceSynchronize() );
#endif

    // Copy result from device memory to host memory
    // h_C contains the result in host memory
    cutilSafeCall( cudaMemcpy(h_C, d_C, c_size, cudaMemcpyDeviceToHost) );
    end=clock();
    duration=double(end-start)/CLOCKS_PER_SEC;
    printf("GPU time: %lf\n",duration);
    // Verify result
    for (i = 0; i < a_row*b_col; ++i) {
        //float sum = h_A[i] + h_B[i];
        if (fabs(h_C[i] - h_C_cpu[i]) > 1e-5)
        {
            //printf("The result is wrong!\n");
            break;
        }
    }

    CleanupResources();
    shrQAFinishExit(argc, (const char **)argv, (i==a_row*b_col) ? QA_PASSED : QA_FAILED);
}
void ZeroInit(float* a, int N)
{
    for(int i=0;i<N;i++)
        a[i]=0;
}
void CleanupResources(void)
{
    // Free device memory
    if (d_A)
        cudaFree(d_A);
    if (d_B)
        cudaFree(d_B);
    if (d_C)
        cudaFree(d_C);

    // Free host memory
    if (h_A)
        free(h_A);
    if (h_B)
        free(h_B);
    if (h_C)
        free(h_C);

    cutilDeviceReset();
}

// Allocates an array with random float entries.
void RandomInit(float* data, int n)
{
    for (int i = 0; i < n; ++i)
        data[i] = rand() / (float)RAND_MAX;
}

// Parse program arguments
void ParseArguments(int argc, char** argv)
{
    for (int i = 0; i < argc; ++i) {
        if (strcmp(argv[i], "--noprompt") == 0 ||
            strcmp(argv[i], "-noprompt") == 0) 
        {
            noprompt = true;
            break;
        }
    }
}

上記は私のCUDAコードです：「MatrixMul.cu」、私のプロジェクトにはこの1つのファイルしかなく、SDK、VectorAddプロジェクトでそれを書き、それを変更するだけです。カーネル関数とメイン関数を 1 つのcuファイルに記述しました。

結果を自分の CPU の結果と比較したところ、同じではないことがわかりました。もう 1 つの問題は、C[rowi*matrixcol+coli] の代わりに tmp 変数を使用したときに、これも間違っていることです。理由もわかりません。

4

1 に答える 1

2

私が見つけたあなたのコードには2つの問題があります。まず第一に、カーネルでは、有効なスレッド インデックスでコードを適切に調整していません。有効なスレッド インデックスは、実際の結果の行列要素に対応するものです。無効なスレッド インデックスは、この領域外のものです。これをチェックしていますが、コード内の間違った場所にあります。これの代わりに：

C[rowi*Bcol+coli]=0; 
for(int k=0;k<Acol;k++) 
{ 
    if(rowi<Arow&&coli<Bcol) 
        C[rowi*Bcol+coli]+=A[rowi*Acol+k]*B[k*Bcol+coli]; 
}

これを使って：

if(rowi<Arow&&coli<Bcol)  {
  C[rowi*Bcol+coli]=0; 
  for(int k=0;k<Acol;k++) 
    { 

        C[rowi*Bcol+coli]+=A[rowi*Acol+k]*B[k*Bcol+coli]; 
    } 
}

コードの記述方法が原因で、有効なスレッド チェックの前にある次のコード行が原因で、有効な範囲外の一部のスレッドが一部の要素をゼロにする必要がありませんでした。

C[rowi*Bcol+coli]=0;

私が見つけた 2 番目の問題は、同等性チェックがおそらく厳しすぎるということです。あなたがこれを持っていた場所：

    if (fabs(h_C[i] - h_C_cpu[i]) > 1e-5)

私はこれを次のように変更しました：

    if (fabs(h_C[i] - h_C_cpu[i]) > 1e-4)

上記の変更により、一致する結果を得ることができました。等価チェックをいじって、一致する桁数を確認できますが、32 ビット浮動小数点数に対して一致する桁数が多すぎると予想していました。ここでの残差チェックはスケーリングされていないため、思ったほどタイトにはなりません。スケーリングされた残差チェックを作成すると、各要素の特定の精度を確実にチェックできます。

さらなる提案として、結果比較ループで、次の行を変更します。

        //printf("The result is wrong!\n");

に：

        printf("The result is wrong at idx: %d CPU: %f GPU: %f\n", i, h_C_cpu[i], h_C[i]);

さらに試してみて、うまくいかない場合に、より有用な結果を得るために。

于 2012-10-07T21:14:53.080 に答える