1

次の (簡略化された) ネストされたループを CUDA 2D カーネルとして移植しようとしています。とのサイズはNgS、NgOデータ セットが大きいほど大きくなります。今のところ、このカーネルがすべての値に対して正しい結果を出力するようにしたいだけです。

// macro that translates 2D [i][j] array indices to 1D flattened array indices
#define idx(i,j,lda) ( (j) + ((i)*(lda)) )

int NgS  = 1859;
int NgO  = 900;

// 1D flattened matrices have been initialized as:
 Radio_cpu = new double [NgS*NgO];
Result_cpu = new double [NgS*NgO];
// ignoring the part where they are filled w/ data

for (m=0; m<NgO; m++) {        
    for (n=0; n<NgS; n++) {
            Result_cpu[idx(n,m,NgO)]] = k0*Radio_cpu[idx(n,m,NgO)]];
    }
}

私が遭遇した例は通常、正方形のループを扱っており、CPU バージョンと比較して、すべての GPU 配列インデックスに対して正しい出力を得ることができませんでした。カーネルを呼び出すホスト コードは次のとおりです。

dim3 dimBlock(16, 16);
dim3 dimGrid;
dimGrid.x = (NgO + dimBlock.x - 1) / dimBlock.x;
dimGrid.y = (NgS + dimBlock.y - 1) / dimBlock.y;

// Result_gpu and Radio_gpu are allocated versions of the CPU variables on GPU
trans<<<dimGrid,dimBlock>>>(NgO, NgS, k0, Radio_gpu, Result_gpu);

カーネルは次のとおりです。

__global__ void trans(int NgO, int NgS,
                      double k0, double * Radio, double * Result) {

int n = blockIdx.x * blockDim.x + threadIdx.x;
int m = blockIdx.y * blockDim.y + threadIdx.y;

if(n > NgS || m > NgO) return;

// map the two 2D indices to a single linear, 1D index
int grid_width = gridDim.x * blockDim.x;
int idxxx = m + (n * grid_width);

Result[idxxx] = k0 * Radio[idxxx];
}

現在のコードを使用して、変数をコピーして戻したResult_cpu変数と比較しました。Result_gpu値を循環すると、次のようになります。

    // matches from NgS = 0...913
    Result_gpu[NgS = 913][NgO = 0]: -56887.2
    Result_cpu[Ngs = 913][NgO = 0]: -56887.2

    // mismatches from NgS = 914...1858
    Result_gpu[NgS = 914][NgO = 0]: -12.2352
    Result_cpu[NgS = 914][NgO = 0]: 79448.6

このパターンは、 の値に関係なく同じですNgO。数時間さまざまな例を見て変更を試すことで、どこで間違いを犯したかを突き止めようとしましたが、これまでのところ、このスキームは目前の明らかな問題を差し引いて機能しましたが、他のスキームはカーネル呼び出しエラーを引き起こしました/左すべての値に対して初期化されていない GPU 配列。間違いがはっきりと見えないので、誰かが私を修正に向けて正しい方向に向けることができれば、本当に感謝しています. 鼻のすぐ下にあり、見えないことは確かです。

問題が発生した場合に備えて、このコードを Kepler カードでテストし、MSVC 2010、CUDA 4.2、および 304.79 ドライバーを使用してコンパイルし、フラグarch=compute_20,code=sm_20とarch=compute_30,code=compute_30フラグの両方を使用してコードをコンパイルしましたが、違いはありません。

4

1 に答える 1

3

@vaca_loca：次のカーネルをテストしました（非正方形のブロック次元でも機能します）：

__global__ void trans(int NgO, int NgS,
                  double k0, double * Radio, double * Result) {

int n = blockIdx.x * blockDim.x + threadIdx.x;
int m = blockIdx.y * blockDim.y + threadIdx.y;
if(n > NgO || m > NgS) return;
int ofs = m * NgO + n;
Result[ofs] = k0 * Radio[ofs];
}

void test() {

int NgS  = 1859, NgO  = 900;
int data_sz = NgS * NgO, bytes = data_sz * sizeof(double);
cudaSetDevice(0);
double *Radio_cpu = new double [data_sz*3],
    *Result_cpu = Radio_cpu + data_sz,
    *Result_gpu = Result_cpu + data_sz;
double k0 = -1.7961233;

srand48(time(NULL));
int i, j, n, m;
for(m=0; m<NgO; m++) {
  for (n=0; n<NgS; n++) {
        Radio_cpu[m + n*NgO] = lrand48() % 234234;
        Result_cpu[m + n*NgO] = k0*Radio_cpu[m + n*NgO];
    }
}

double *g_Radio, *g_Result;
cudaMalloc((void **)&g_Radio, bytes * 2);
g_Result = g_Radio + data_sz;
cudaMemcpy(g_Radio, Radio_cpu, bytes, cudaMemcpyHostToDevice);

dim3 dimBlock(16, 16);
dim3 dimGrid;
dimGrid.x = (NgO + dimBlock.x - 1) / dimBlock.x;
dimGrid.y = (NgS + dimBlock.y - 1) / dimBlock.y;

trans<<<dimGrid,dimBlock>>>(NgO, NgS, k0, g_Radio, g_Result);

cudaMemcpy(Result_gpu, g_Result, bytes, cudaMemcpyDeviceToHost);

for(m=0; m<NgO; m++) {
    for (n=0; n<NgS; n++) {
        double c1 = Result_cpu[m + n*NgO],
                c2 = Result_gpu[m + n*NgO];
        if(std::abs(c1-c2) > 1e-4)
            printf("(%d;%d): %.7f %.7f\n", n, m, c1, c2);
    }
}
cudaFree(g_Radio);
delete []Radio_cpu;
}

ただし、私の意見では、クワッドを使用してグローバルメモリからデータにアクセスすることは、アクセスストライドがかなり大きいため、キャッシュにあまり適していません。アルゴリズムが2Dローカリティのデータにアクセスすることが重要な場合は、代わりに2Dテクスチャの使用を検討してください。

于 2012-08-05T11:09:18.720 に答える