CUBLAS的同步

Question

CUBLAS文档提到，在读取标量结果之前我们需要同步：

“另外，返回标量结果的几个函数，如amax()、amin、asum()、rotg()、rotmg()、dot()和nrm2()，通过在主机或设备上引用返回结果值。注意，即使这些函数立即返回，类似于矩阵和向量结果，标量结果也只有在GPU上执行例程时才能准备好。这需要适当的同步，以便从主机读取结果。”

这是否意味着我们应该始终在读取主机的标量结果之前进行同步，即使我们只使用单个流？我一直在寻找一个关于NVIDIA的CUDA文档的例子，但没有找到一个。

但在NVIDIA提供的共轭梯度示例中，有以下代码

while (r1 > tol*tol && k <= max_iter)
{
    if (k > 1)
    {
        b = r1 / r0;
        cublasStatus = cublasSscal(cublasHandle, N, &b, d_p, 1);
        cublasStatus = cublasSaxpy(cublasHandle, N, &alpha, d_r, 1, d_p, 1);
    }
    else
    {
        cublasStatus = cublasScopy(cublasHandle, N, d_r, 1, d_p, 1);
    }

    cusparseScsrmv(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &alpha, descr, d_val, d_row, d_col, d_p, &beta, d_Ax);
    cublasStatus = cublasSdot(cublasHandle, N, d_p, 1, d_Ax, 1, &dot);
    a = r1 / dot;

    cublasStatus = cublasSaxpy(cublasHandle, N, &a, d_p, 1, d_x, 1);
    na = -a;
    cublasStatus = cublasSaxpy(cublasHandle, N, &na, d_Ax, 1, d_r, 1);

    r0 = r1;
    cublasStatus = cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
    cudaThreadSynchronize();
    printf("iteration = %3d, residual = %e\n", k, sqrt(r1));
    k++;
}

这里有一个cudaThreadSynchronize()调用，在while循环结束之前。是给cublasSdot电话的吗？但是循环中有两个cublasSdot调用。为什么在第二个cudaThreadSynchronize之后有一个cublasSdot ()，而不是第一个cublasSdot呢？

编辑:为了了解发生了什么，我使用了以下代码来比较有同步和不同步的点积结果。

int main(int argc, char **argv)
{
    /* Generate a large vector */
    int N = 1024 * 1024 * 512;

    double *x_cpu = (double *)malloc(sizeof(double)*N);
    for (int i = 0; i < N; i++)
    {
        x_cpu[i] = double(rand()) / RAND_MAX;
    }


    double *x_gpu;
    cudaMalloc((void **)&x_gpu, N*sizeof(double));
    cudaMemcpy(x_gpu, x_cpu, N*sizeof(double), cudaMemcpyHostToDevice);

    /* Get handle to the CUBLAS context */
    cublasHandle_t cublasHandle = 0;
    cublasStatus_t cublasStatus;
    cublasStatus = cublasCreate(&cublasHandle);

    int M = 1000;
    std::vector<double> x_dot_vec(M, 0.0);
    double *x_dot_ptr = &(x_dot_vec[0]);

    std::cout << "Begin Launching CUBLAS........" << std::endl;

    for(int j = 0; j < M; j++){
        cublasDdot(cublasHandle, N, x_gpu, 1, x_gpu, 1, x_dot_ptr + j);
    }

    std::cout << "End Launching CUBLAS........." << std::endl;

    double old_value = x_dot_vec.back();
    cudaDeviceSynchronize();
    double new_value = x_dot_vec.back();
    std::cout << "Old Value: " << old_value << ",   New Value: " << new_value << std::endl;

    free(x_cpu);
    cudaFree(x_gpu);

    return 0;
}

这里的想法是，我们创建一个非常大的向量，并多次使用cublas计算它的点积，并将返回值写入主机上的数组中。在启动所有cublas函数之后，我们读取结果数组的最后一个元素而不进行同步。如果cublasDdot调用实际上是非阻塞的，那么还不应该编写最后一个元素。然后我们进行同步，并再次读取最后一个元素。这一次，它应该存储正确的点积，希望给我们一个不同的价值，我们得到了一个没有同步。但是，当我运行这段代码时，这两个值总是相同的。在cublas调用之前和之后，输出之间需要很长时间。看起来cublasDdot实际上是阻塞的，不像在CUBLAS文档中所说的那样。

我还尝试了下面的版本，其中的结果输出到一个设备数组，而不是主机数组。但是结果看起来是一样的。

int main(int argc, char **argv)
{
    /* Generate a large vector */
    int N = 1024 * 1024 * 512;

    double *x_cpu = (double *)malloc(sizeof(double)*N);
    for (int i = 0; i < N; i++)
    {
        x_cpu[i] = double(rand()) / RAND_MAX;
    }


    double *x_gpu;
    cudaMalloc((void **)&x_gpu, N*sizeof(double));
    cudaMemcpy(x_gpu, x_cpu, N*sizeof(double), cudaMemcpyHostToDevice);

    /* Get handle to the CUBLAS context */
    cublasHandle_t cublasHandle = 0;
    cublasStatus_t cublasStatus;
    cublasStatus = cublasCreate(&cublasHandle);
    cublasSetPointerMode(cublasHandle, CUBLAS_POINTER_MODE_DEVICE);

    int M = 1000;
    std::vector<double> x_dot_vec(M, 0.0);
    double *x_dot_ptr = &(x_dot_vec[0]);
    double *dot_gpu;
    cudaMalloc((void **)&dot_gpu, sizeof(double) * M);
    cudaMemcpy(dot_gpu, x_dot_ptr, M * sizeof(double), cudaMemcpyHostToDevice);
    double old_value, new_value;

    std::cout << "Begin Launching CUBLAS........" << std::endl;

    for(int j = 0; j < M; j++){
        cublasDdot(cublasHandle, N, x_gpu, 1, x_gpu, 1, dot_gpu + j);
    }

    std::cout << "End Launching CUBLAS........." << std::endl;

    cudaMemcpy(&old_value, dot_gpu + M - 1, sizeof(double), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    cudaMemcpy(&new_value, dot_gpu + M - 1, sizeof(double), cudaMemcpyDeviceToHost);
    std::cout << "Old Value: " << old_value << ",   New Value: " << new_value << std::endl;

    free(x_cpu);
    cudaFree(x_gpu);
    cudaFree(dot_gpu);

    return 0;
}

Answer 1

我认为那个代码是不正确的。正如您已经注意到的，在CUBLAS V2 API中，cublasSdot是一个非阻塞调用，理论上需要一个同步原语才能在主机内存中使用该结果。

第一个cublasSdot调用还应该有一个同步点，类似于：

...
cusparseScsrmv(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &alpha, descr, d_val, d_row, d_col, d_p, &beta, d_Ax);
cublasStatus = cublasSdot(cublasHandle, N, d_p, 1, d_Ax, 1, &dot);
cudaDeviceSynchronize();
a = r1 / dot;
...

该示例代码还使用了长时间不推荐的cudaThreadSynchronize API调用。我建议向NVIDIA公司提交一份关于这两个项目的错误报告。