c - 尝试消除共轭梯度内核中的 cudaDeviceSynchronize()

Question

我正在具有统一内存的 TegraK1 板上实现共轭梯度求解器。我面临的问题是在循环中我必须执行 cudaDeviceSynchronize();两次更新变量，与 TI Keystone-II 相比，这对我的性能造成了很多倍的影响，虽然 TI Keystone-II 的计算能力较低，但我只使用 Naive 代码来并行计算数据。我正在使用 CUDA 版本 - 6.0。

....    
double *A, *b, *x, *x1, *r, *d, *q, deltaNew, deltaFirst, Alpha, deltaOld, Beta; // data for init processing

double *temp, *temp1, Alpha1;

b = (double *) malloc(sizeof(double)*N*1); // b original
x1 = (double *) malloc(sizeof(double)*N*1); // x1 

checkCudaErrors(cudaMallocManaged(&A, sizeof(double)*N*N)); // A original 
checkCudaErrors(cudaMallocManaged(&x, sizeof(double)*N*1)); // x original
checkCudaErrors(cudaMallocManaged(&r, sizeof(double)*N*1)); // r original 
checkCudaErrors(cudaMallocManaged(&d, sizeof(double)*N*1)); // d original
checkCudaErrors(cudaMallocManaged(&q, sizeof(double)*N*1)); // q original
checkCudaErrors(cudaMallocManaged(&temp, sizeof(double)*1*1)); //  temp of d'*q for temporary storage
checkCudaErrors(cudaMallocManaged(&temp1, sizeof(double)*1*1)); //  temp1 of r'*r for temporary storage

fprintf(stderr, "\nIntializing data\n");
// Intializing all the data
setup_data(&A[0], &b[0], &x[0], &r[0], &d[0], &deltaNew, &deltaFirst);

// Get handle to the CUBLAS context 
cublasHandle_t cublasHandle = 0;
cublasCreate(&cublasHandle);
fprintf(stderr, "\nData setup done.. Starting..\n");    
startTime_GPU = omp_get_wtime();

while(deltaNew > (EPSI)*deltaFirst)
{
    // cublasSgemm(handle, op, op, colof2, rowof1, colof1, scalar1, mat2, colof2, mat1, colof1, scalar2, result, colof2 );
    cublasDgemm(cublasHandle, CUBLAS_OP_N, CUBLAS_OP_N, 1, N, N, &alpha, d, 1, A, N, &beta, q, 1);        // q = A * d   
    cublasDgemm(cublasHandle, CUBLAS_OP_N, CUBLAS_OP_N, 1, 1, N, &alpha, q, 1, d, N, &beta, temp, 1);        // alpha = deltaNew/(d' * q);

    cudaDeviceSynchronize(); // POSSIBLY ELIMINATE THIS
    Alpha = deltaNew/temp[0]; // alpha = deltaNew/(d' * q);
    Alpha1 = (-1)*Alpha;   

    // cublasSaxpy(handle, N, scalar, scaledinput, stride1, inout, stride2);
    cublasDaxpy(cublasHandle, N, &Alpha, d, 1, x, 1); // x = x + alpha * d  
    cublasDaxpy(cublasHandle, N, &Alpha1, q, 1, r, 1); // r = r - alpha * q

    deltaOld = deltaNew; // deltaOld = deltaNew  

    cublasDgemm(cublasHandle, CUBLAS_OP_N, CUBLAS_OP_N, 1, 1, N, &alpha, r, 1, r, N, &beta, temp1, 1);        // deltaNew = r' * r

    cudaDeviceSynchronize(); // POSSIBLY ELIMINATE THIS
    deltaNew = temp1[0];
    Beta = deltaNew/deltaOld; // beta = deltaNew/deltaOld
    cublasDgeam(cublasHandle, CUBLAS_OP_N, CUBLAS_OP_N, 1, N, &alpha, r, 1, &Beta, d, 1, d, 1); // d = r + beta * d
}

endTime_GPU = omp_get_wtime();
...
...

任何人都可以提出任何改进或增强建议，我可以用它来消除或减少这些改进或增强，从而提高性能。

Answer 1

在 CUDA 中，内核启动始终是异步的(即使在 CUDA 1.0 中也是如此)。当时，为了让 CUDA 内核运行，需要在设备内存之间进行显式 memcpy。 CPU/GPU 同步被隐藏，因为设备->主机 memcpy 会隐式强制执行顺序:设备->主机 memcpy 在内核完成之前无法启动。

当CPU和GPU映射相同的内存时，必须更明确地进行同步。如果没有 cudaDeviceSynchronize()，您的 CPU 代码可能会在 GPU 写入 DGEMM 之前读取输出。托管内存在内核启动和 CPU/GPU 同步事件(例如 cudaDeviceSynchronize())背后实现了许多隐式策略，以简化编程。

获得 CPU/GPU 并发性并仍然正确同步的方法是通过多缓冲并将 CUDA 事件与每个缓冲区关联。在每个 DGEMM 之后调用 cudaEventRecord()，并在使用结果之前对该事件调用 cudaEventWait()。

talonmies 已经说过了，但值得重复:如果你想要良好的性能，你可能不得不放弃托管内存。