试图消除共轭梯度核中的 cudaDeviceSynchronize()
Trying to eliminate cudaDeviceSynchronize() in conjugate gradient kernel
我正在具有统一内存的 TegraK1 板上实现共轭梯度求解器。我面临的问题是在循环中我必须做 cudaDeviceSynchronize();两次更新变量,与 TI Keystone-II 相比,这对我的性能造成了很多倍的伤害,TI Keystone-II 虽然计算能力较低,但我只使用朴素代码来并行计算数据。我正在使用 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();
...
...
任何人都可以提出任何改进或增强的建议,我可以使用这些改进或增强来消除或减少这些改进或增强,从而提高性能。
内核启动在 CUDA 中总是是异步的(即使在 CUDA 1.0 中)。那时,需要显式 memcpy 进出设备内存,以便 CUDA 内核运行。 CPU/GPU 同步被隐藏了,因为设备->主机 memcpy 会隐含地强制执行一个顺序:设备->主机 memcpy 在内核完成之前无法启动。
当CPU和GPU映射相同的内存时,同步必须更明确地完成。如果没有您的 cudaDeviceSynchronize()
,您的 CPU 代码可能会在 GPU 写入之前从 DGEMM 读取输出。托管内存在内核启动和 CPU/GPU 同步事件(例如 cudaDeviceSynchronize()
之后实施了很多隐式策略,以使编程更容易。
获得 CPU/GPU 并发性并仍然正确同步的方法是通过多缓冲并将 CUDA 事件与每个缓冲区相关联。在每个 DGEMM 之后调用 cudaEventRecord()
,并在使用结果之前对该事件调用 cudaEventWait()
。
talonmies 已经说过,但值得重复:如果你想要好的性能,你可能不得不放弃托管内存。
我正在具有统一内存的 TegraK1 板上实现共轭梯度求解器。我面临的问题是在循环中我必须做 cudaDeviceSynchronize();两次更新变量,与 TI Keystone-II 相比,这对我的性能造成了很多倍的伤害,TI Keystone-II 虽然计算能力较低,但我只使用朴素代码来并行计算数据。我正在使用 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();
...
...
任何人都可以提出任何改进或增强的建议,我可以使用这些改进或增强来消除或减少这些改进或增强,从而提高性能。
内核启动在 CUDA 中总是是异步的(即使在 CUDA 1.0 中)。那时,需要显式 memcpy 进出设备内存,以便 CUDA 内核运行。 CPU/GPU 同步被隐藏了,因为设备->主机 memcpy 会隐含地强制执行一个顺序:设备->主机 memcpy 在内核完成之前无法启动。
当CPU和GPU映射相同的内存时,同步必须更明确地完成。如果没有您的 cudaDeviceSynchronize()
,您的 CPU 代码可能会在 GPU 写入之前从 DGEMM 读取输出。托管内存在内核启动和 CPU/GPU 同步事件(例如 cudaDeviceSynchronize()
之后实施了很多隐式策略,以使编程更容易。
获得 CPU/GPU 并发性并仍然正确同步的方法是通过多缓冲并将 CUDA 事件与每个缓冲区相关联。在每个 DGEMM 之后调用 cudaEventRecord()
,并在使用结果之前对该事件调用 cudaEventWait()
。
talonmies 已经说过,但值得重复:如果你想要好的性能,你可能不得不放弃托管内存。