cuda - マルチ GPU cuda プログラミング: 複数のデバイスから 1 つのカーネルを起動する

Question

複数の GPU で cuda を使用してコーディングする方法を学習しています。私のデバイスの計算能力は 4.0 なので、1 つのホスト スレッドを使用して複数の GPU の cuda をコーディングできることを理解しています。まず、SDK のサンプルの 1 つ「simpleMultiGPU.cpp」を参照し、1 つまたは 2 つの GPU を使用して実行しました。どちらの場合もパフォーマンスは非常に似ていることがわかりました。つまり、2 つの GPU でコードを実行すると、1 つの GPU で実行するのと同じくらい遅い/速いということです。最適化されたコードではないことは理解していますが、これは実際に私の目の前にある例であり、機能していません。2 つのデバイスのコードはシリアルで実行されていると思います。

この SDK の例を実行しているときに、誰かが同じ問題を経験しましたか?

この SDK の例に基づいて単純なベクトル加算コードを作成しましたが、同様の方法で実行されます (予想どおり)。非同期の cuda 呼び出しを使用しており、固定されたホスト メモリも使用しています。この動作の背後にある理由を理解しようとしています。

どんな洞察も高く評価されます。

メインコードのコピーは次のとおりです。

typedef struct {
float* vec;
int N;
} vector;

extern "C" {

//Define kernel for vector addition
__global__ void vecadd_kernel(float *avec, int N, float* bvec, float *cvec){
int tId=blockIdx.x*blockDim.x+threadIdx.x;

if(tId < N)
    cvec[tId]=avec[tId]+bvec[tId];


}

void launch_addvec_kernel(float *avec, int N, float* bvec, float *cvec, int THREAD_N, int BLOCK_N, cudaStream_t &s){

vecadd_kernel<<< BLOCK_N, THREAD_N, 0, s >>> (avec,N,bvec,cvec);
        getLastCudaError("reduceKernel() execution failed.\n");

}


}

int main(){

clock_t lapse;
float cpu_time;
lapse=clock();
vector avec, bvec, cvec, cvec_gpu;

int N=256*256*256;
int threads=256;

avec.N=N; 
bvec.N=avec.N; 
cvec.N=avec.N; 

avec.vec=(float*)malloc(sizeof(float)*avec.N);
bvec.vec=(float*)malloc(sizeof(float)*bvec.N);
cvec.vec=(float*)malloc(sizeof(float)*cvec.N);
cvec_gpu.vec=(float*)malloc(sizeof(float)*avec.N);

for(int i=0;i<avec.N;++i){
    avec.vec[i]=i;
    bvec.vec[i]=i;
}


//Normal CPU addition
#pragma unroll
for(int i=0;i<avec.N;++i){
    cvec.vec[i]=avec.vec[i]+bvec.vec[i];
}

cpu_time=clock()-lapse;

printf("CPU execution time = %f seconds \n",cpu_time/CLOCKS_PER_SEC);

//-------------------- Multi-GPU code -------------------------------
//-------------------- Multi-GPU code -------------------------------
//-------------------- Multi-GPU code -------------------------------

//Get number of CUDA enabled devices
lapse=clock();
int deviceCount;
cudaGetDeviceCount(&deviceCount);
//deviceCount=1;

vector apartvecs[deviceCount], bpartvecs[deviceCount], cpartvecs[deviceCount];
vector apartvecs_gpu[deviceCount], bpartvecs_gpu[deviceCount],  cpartvecs_gpu[deviceCount];

int i,j;

    //Subdividing input data across GPUs
    //Get data sizes for each GPU
for (i=0; i<deviceCount; ++i)
    apartvecs[i].N = N/deviceCount;

    //Take into account "odd" data sizes
for (i=0; i<N%deviceCount; ++i)
    ++apartvecs[i].N; 


int offset[deviceCount];

offset[0]=0;
offset[1]=apartvecs[0].N;

cudaStream_t stream[deviceCount];

    //Create streams for issuing GPU command asynchronously and allocate memory (GPU and System page-locked)
for (i=0; i<deviceCount; ++i){
        checkCudaErrors( cudaSetDevice(i) );
        checkCudaErrors( cudaStreamCreate(&stream[i]) );

    cpartvecs[i].vec=(float*)malloc(sizeof(float)*apartvecs[i].N);

    memset(cpartvecs[i].vec,'\0',sizeof(float)*apartvecs[i].N);

    //Allocate device memory
        checkCudaErrors( cudaMalloc((void**)&apartvecs_gpu[i].vec, apartvecs[i].N * sizeof(float)) );
        checkCudaErrors( cudaMalloc((void**)&bpartvecs_gpu[i].vec, apartvecs[i].N * sizeof(float)) );
        checkCudaErrors( cudaMalloc((void**)&cpartvecs_gpu[i].vec, apartvecs[i].N * sizeof(float)) );

    //Allocate pinned memory on host
        checkCudaErrors( cudaMallocHost((void**)&apartvecs[i].vec, apartvecs[i].N * sizeof(float)));
        checkCudaErrors( cudaMallocHost((void**)&bpartvecs[i].vec, apartvecs[i].N * sizeof(float)));

    for (j=0;j<apartvecs[i].N;++j){
        int j1=j+offset[i]; 
        apartvecs[i].vec[j]=avec.vec[j1];
        bpartvecs[i].vec[j]=bvec.vec[j1];
        //printf("%d \t %d \t %d \t %d \t %f\n",i,j,offset[i],j1,apartvecs[i].vec[j]);
    }
}

    //Copy data to GPU, launch the kernel and copy data back. All asynchronously
for (i=0; i<deviceCount; ++i){

        //Set device
        checkCudaErrors( cudaSetDevice(i) );

        //Copy input data from CPU
        checkCudaErrors( cudaMemcpyAsync(apartvecs_gpu[i].vec, apartvecs[i].vec, apartvecs[i].N * sizeof(float), cudaMemcpyHostToDevice, stream[i]) );
        checkCudaErrors( cudaMemcpyAsync(bpartvecs_gpu[i].vec, bpartvecs[i].vec, apartvecs[i].N * sizeof(float), cudaMemcpyHostToDevice, stream[i]) );

    int numblocks = N/threads;

//  printf("before kernel %d \n",apartvecs[i].N);
    launch_addvec_kernel(apartvecs_gpu[i].vec,apartvecs[i].N,bpartvecs_gpu[i].vec,cpartvecs_gpu[i].vec,threads,numblocks,stream[i]);

        //Read back GPU results

        checkCudaErrors( cudaMemcpyAsync(cpartvecs[i].vec, cpartvecs_gpu[i].vec, apartvecs[i].N * sizeof(float), cudaMemcpyDeviceToHost, stream[i]) );

    //printf("here 5\n");

}

    //Process GPU results
    for(i = 0; i < deviceCount; i++){

        //Set device
        checkCudaErrors( cudaSetDevice(i) );

        //Wait for all operations to finish
        cudaStreamSynchronize(stream[i]);

// cudaDeviceSynchronize();

    for(int j=0; j<apartvecs[i].N; ++j){
        int j1=j+offset[i]; 
        cvec_gpu.vec[j1]=cpartvecs[i].vec[j];
        //printf("%d \t %d \t %d \t %d \t %f\n",i,j,offset[i],j1,cvec_gpu.vec[j1]);
    }

        //Shut down this GPU
        checkCudaErrors( cudaFreeHost(apartvecs[i].vec) );
        checkCudaErrors( cudaFreeHost(bpartvecs[i].vec) );
        checkCudaErrors( cudaFree(apartvecs_gpu[i].vec) );
        checkCudaErrors( cudaFree(bpartvecs_gpu[i].vec) );
        checkCudaErrors( cudaFree(cpartvecs_gpu[i].vec) );
        checkCudaErrors( cudaStreamDestroy(stream[i]) );

}

free(avec.vec);
free(bvec.vec);
free(cvec.vec);
free(cvec_gpu.vec);
cpu_time=clock()-lapse;

printf("GPU execution time = %f seconds \n",cpu_time/CLOCKS_PER_SEC);

}