0

我正在学习如何在多个 GPU 上使用 cuda 进行编码。我的设备的计算能力是 4.0，所以我知道我可以使用一个主机线程为多个 GPU 编写 cuda。首先，我参考了 SDK 示例之一“simpleMultiGPU.cpp”并使用一个和两个 GPU 运行它。我发现两种情况下的性能非常相似，即在两个 GPU 中运行代码与在一个 GPU 中运行代码一样慢/快。我知道它不是一个优化的代码，但这确实是我面前的一个例子，这是行不通的。我认为这两个设备中的代码是串行运行的。

有人在运行此 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);

}

4

0 回答 0