如何减少 CudaMemcpy 开销

Question

我有一个大小为 3000 的数组，该数组包含 0 和 1。我想找到从第 0 个索引开始在该位置存储 1 的第一个数组位置。我将此数组传输到主机，并在设备上计算此数组。然后我按顺序计算主机上的索引。在我的程序中，我想重复执行此计算 4000 次或更多次。我想减少此过程所花费的时间。有没有其他方法可以做到这一点，并计算这个数组实际上在 GPU 上，所以我每次都必须传输它。

int main()
{
for(int i=0;i<4000;i++)
{
    cudaMemcpy(A,dev_A,sizeof(int)*3000,cudaMemcpyDeviceToHost);
    int k;
    for(k=0;k<3000;k++)
    {
        if(A[k]==1)
        {
            break;
        } 
    }
    printf("got k is %d",k);
}
} 

完整的代码是这样的 #include“cuda.h” ＃包括 #定义尺寸 2688 #define 块 14 #define 线程 192

__global__ void kernel(int *A,int *d_pos)
{
int thread_id=threadIdx.x+blockIdx.x*blockDim.x;
while(thread_id<SIZE)
{
    if(A[thread_id]==INT_MIN)
    {
        *d_pos=thread_id;
        return;
    }
    thread_id+=1;   
}

}

__global__ void kernel1(int *A,int *d_pos)
{
int thread_id=threadIdx.x+blockIdx.x*blockDim.x;
if(A[thread_id]==INT_MIN)
{
    atomicMin(d_pos,thread_id);
}

}

int main()
{
int pos=INT_MAX,i;
int *d_pos;
int A[SIZE];
int *d_A;
for(i=0;i<SIZE;i++)
{
    A[i]=78;
}
A[SIZE-1]=INT_MIN;
cudaMalloc((void**)&d_pos,sizeof(int));
cudaMemcpy(d_pos,&pos,sizeof(int),cudaMemcpyHostToDevice);
cudaMalloc((void**)&d_A,sizeof(int)*SIZE);
cudaMemcpy(d_A,A,sizeof(int)*SIZE,cudaMemcpyHostToDevice);

cudaEvent_t start_cp1,stop_cp1;
    cudaEventCreate(&stop_cp1);
    cudaEventCreate(&start_cp1);
    cudaEventRecord(start_cp1,0);

kernel1<<<BLOCKS,THREADS>>>(d_A,d_pos);

cudaEventRecord(stop_cp1,0);
    cudaEventSynchronize(stop_cp1);
    float elapsedTime_cp1;
    cudaEventElapsedTime(&elapsedTime_cp1,start_cp1,stop_cp1);
    cudaEventDestroy(start_cp1);
    cudaEventDestroy(stop_cp1);
    printf("\nTime taken by kernel is  %f\n",elapsedTime_cp1);
cudaDeviceSynchronize();

cudaEvent_t start_cp,stop_cp;
    cudaEventCreate(&stop_cp);
    cudaEventCreate(&start_cp);
    cudaEventRecord(start_cp,0);

cudaMemcpy(A,d_A,sizeof(int)*SIZE,cudaMemcpyDeviceToHost);

    cudaEventRecord(stop_cp,0);
    cudaEventSynchronize(stop_cp);
    float elapsedTime_cp;
    cudaEventElapsedTime(&elapsedTime_cp,start_cp,stop_cp);
    cudaEventDestroy(start_cp);
    cudaEventDestroy(stop_cp);
    printf("\ntime taken by copy of an array is  %f\n",elapsedTime_cp);






    cudaEvent_t start_cp2,stop_cp2;
    cudaEventCreate(&stop_cp2);
    cudaEventCreate(&start_cp2);
    cudaEventRecord(start_cp2,0);

    cudaMemcpy(&pos,d_pos,sizeof(int),cudaMemcpyDeviceToHost);

    cudaEventRecord(stop_cp2,0);
    cudaEventSynchronize(stop_cp2);
    float elapsedTime_cp2;
    cudaEventElapsedTime(&elapsedTime_cp2,start_cp2,stop_cp2);
    cudaEventDestroy(start_cp2);
    cudaEventDestroy(stop_cp2);
    printf("\ntime taken by copy of a variable is  %f\n",elapsedTime_cp2);


cudaMemcpy(&pos,d_pos,sizeof(int),cudaMemcpyDeviceToHost);
printf("\nminimum index is %d\n",pos);
return 0;
}

如何使用任何其他性能选项来减少此代码所花费的总时间。

Answer 1

如果您在 GPU 上运行内核 4000 次，则可能需要通过不同的流在内核上使用异步执行。使用 cudaMemCpyAsync 可能会更快，它是主机的非阻塞函数（在您执行 M 次内核的情况下）。

流和异步执行的快速介绍： https://devblogs.nvidia.com/parallelforall/how-overlap-data-transfers-cuda-cc/

流和并发： http://on-demand.gputechconf.com/gtc-express/2011/presentations/StreamsAndConcurrencyWebinar.pdf

希望这可以帮助...