1

ここ数か月、並列プログラミングについて勉強してきましたが、今は自分のアプリケーションをマルチ GPU プラットフォームに適応させようとしています。問題は、複数の GPU を使用して配列を反復処理する方法をまだよく理解していないことです。

メイン配列を小さなサブ配列に分割し、それぞれを各 GPU に送信する必要がありますか、または配列のフラグメントで各 GPU を反復させる方法がありますか? このアプリケーションのシリアルおよびシングル GPU バージョンが動作しており、さまざまな方法を使用してこの問題を解決し、マルチ GPU に適応させようとしましたが、いずれも以前の 2 つのバージョンと同じ結果を返しません。これ以上何ができるかわからないので、マルチ GPU システムで配列を反復処理する方法を理解していないというのが私の結論です。誰か助けてくれませんか？

私のコードは N 回の反復を実行し、各反復で (グリッドを表す) 配列の各値を調べて、新しい値を計算します。

これは、私のコードが現在どのように見えるかのスケッチです。

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>

#define DIM     24
#define BLOCK_SIZE 16
#define SRAND_VALUE 585

__global__ void random(int* t, int* newT){

    int iy = blockDim.y * blockIdx.y + threadIdx.y + 1;
    int ix = blockDim.x * blockIdx.x + threadIdx.x + 1;
    int id = iy * (dim+2) + ix;

    if (iy <= DIM && ix <= DIM) {
        if (t[id] % 2 == 0)
            newT[id] = t[id]*3;
        else
            newT[id] = t[id]*5;
    }
}

int main(int argc, char* argv[]){
    int i,j, devCount;
    int *h_test, *d_test, *d_tempTest, *d_newTest;
    size_t gridBytes;

    cudaGetDeviceCount(&devCount);

    gridBytes = sizeof(int)*(DIM)*(DIM);
    h_test = (int*)malloc(gridBytes);

    srand(SRAND_VALUE);
    #pragma omp parallel for private(i,j)
        for(i = 1; i<=DIM;i++) {
            for(j = 1; j<=DIM; j++) {
                h_test[i*(DIM)+j] = rand() % 2;
            }
        }

    if (devCount == 0){
        printf("There are no devices in this machine!");
        return 1; // if there is no GPU, then break the code
    }

    dim3 blockSize(BLOCK_SIZE, BLOCK_SIZE,1);
    int  linGrid = (int)ceil(DIM/(float)BLOCK_SIZE);
    dim3 gridSize(linGrid,linGrid,1);

    dim3 cpyBlockSize(BLOCK_SIZE,1,1);
    dim3 cpyGridRowsGridSize((int)ceil(DIM/(float)cpyBlockSize.x),1,1);
    dim3 cpyGridColsGridSize((int)ceil((DIM+2)/(float)cpyBlockSize.x),1,1);

    else if (devCount == 1){

        cudaMalloc(&d_test, gridBytes);
        cudaMalloc(&d_tempTest, gridBytes);
        cudaMalloc(&d_newTest, gridBytes);

        cudaMemcpy(d_test, h_test, gridBytes, cudaMemcpyHostToDevice);

        for (iter = 0; iter < DIM; iter ++){
            random<<<gridSize, blockSize>>>(d_test, d_newTest);

            d_tempTest = d_test;
            d_test = d_newTest;
            d_newTest = d_tempTest;
        }

        cudaMemcpy(h_test, d_test, gridBytes, cudaMemcpyDeviceToHost);

        return 0;
    }

    else{
        int nThreads, tId, current;
        omp_set_num_threads(devCount);

        for (iter = 0; iter < DIM; iter ++){

            #pragma omp parallel private(tId, h_subGrid, ) shared(h_grid, gridBytes)
            {
                tId = omp_get_thread_num();
                cudaSetDevice(tId);

                cudaMalloc(&d_test, gridBytes);
                cudaMalloc(&d_tempTest, gridBytes);
                cudaMalloc(&d_newTest, gridBytes);

                cudaMemcpy(d_grid, h_grid, gridBytes, cudaMemcpyHostToDevice);

                ******// What do I do here//******

            } 
        }
        return 0;
    }
}

前もって感謝します。

4

1 に答える 1

1

簡単な答え: はい、GPU ごとに配列をサブ配列に分割する必要があります。

詳細: 各 GPU には独自のメモリがあります。コードでは、各 GPU で配列全体にメモリを割り当て、配列全体を各 GPU にコピーします。これで、配列のサブセットを操作できるようになりました。ただし、コピーバックする場合は、各配列の更新された部分のみをコピーするようにする必要があります。最初からより良い方法は、特定の GPU で更新したい配列の部分だけをコピーすることです。

解決策: multiGPU 部分を次のように変更します (要素を見逃さないようにする必要がありますgridBytes%devCount != 0。私のコード スニペットはこれをチェックしません)。

int gridBytesPerGPU = gridBytes/devCount;
cudaMalloc(&d_test, gridBytesPerGPU);
cudaMalloc(&d_newTest, gridBytesPerGPU );

cudaMemcpy(d_test, &h_test[tId*gridBytesPerGPU], gridBytesPerGPU, cudaMemcpyHostToDevice); // copy only the part of the array that you want to use on that GPU
// do the calculation
cudaMemcpy(&h_test[tId*gridBytesPerGPU], d_newTest, gridBytesPerGPU, cudaMemcpyDeviceToHost);

これで、適切なブロックとグリッド サイズを計算するだけで済みます。以下の (c) を参照してください。その部分に問題がある場合は、コメントで質問してください。この回答を延長します。

それとは別に、あなたのコードには私が理解できない部分がいくつかあります:

a) なぜポインターを交換する必要があるのですか?

b) カーネル部分を複数回実行しますが、for ループ内のコードはカウンターに依存しません。なんで？何が恋しいですか？

for (iter = 0; iter < DIM; iter ++){
    random<<<gridSize, blockSize>>>(d_test, d_newTest);

    d_tempTest = d_test;
    d_test = d_newTest;
    d_newTest = d_tempTest;
}

c) この単純なカーネルのグリッドとブロック サイズの計算は少し複雑に見えます (質問を読んだときにスキップしました)。問題を 1 次元の問題と見なすと、カーネルを含めてすべてがはるかに単純に見えます。

于 2015-07-16T20:57:44.840 に答える