cuda - CUBLAS 同步

Question

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

“此外，少数返回标量结果的函数，例如 amax()、amin、asum()、rotg()、rotmg()、dot() 和 nrm2()，通过引用返回结果值主机或设备。请注意，即使这些函数立即返回，与矩阵和向量结果类似，标量结果仅在 GPU 上的例程执行完成时才准备好。这需要适当的同步才能从主机。”

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

但是在 conjugate gradient example由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++;
}

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

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

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 提交有关这两项的错误报告。