试图消除共轭梯度内核中的 cudaDeviceSynchronize()答案

【问题标题】：Trying to eliminate cudaDeviceSynchronize() in conjugate gradient kernel试图消除共轭梯度内核中的 cudaDeviceSynchronize()
【发布时间】：2015-07-06 18:25:32
【问题描述】：

我在具有统一内存的 TegraK1 板上实现共轭梯度求解器。我面临的问题是在循环中我必须做 cudaDeviceSynchronize();两次更新变量，与 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();
...
...

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

【问题讨论】：

最明显的优化是首先不使用托管内存。这可能是潜在的性能问题。
我认为我必须使用托管内存，因为它是统一内存并且使用普通的 malloc 不起作用。
显然正常的 malloc 不起作用。但是，如果您在设备内存中显式分配并保留大部分矩阵和向量，并显式管理从主机到设备的传输，您可能会为性能提供更好的资金
你怎么在共轭梯度法中有3个矩阵-矩阵乘法（cublasDgemm）？另外我猜你想要矩阵向量乘法。

标签： c cuda parallel-processing gpu gpgpu

【解决方案1】：

内核启动始终在 CUDA 中是异步的（即使在 CUDA 1.0 中也是如此）。当时，需要显式的 memcpy 进出设备内存才能使 CUDA 内核运行。 CPU/GPU 同步被隐藏了，因为 device->host memcpy 将隐式强制执行排序：device->host memcpy 直到内核完成后才能启动。

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

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

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

【讨论】：