如何正确测量 CUDA 时间？

Question

我试图正确测量并行和顺序执行的时间，但我对此表示怀疑：

假设我们有以下代码：

    //get the time
    clock_t start,finish;
    double totaltime;
    start = clock(); 

    double *d_A, *d_B, *d_X;

    cudaMalloc((void**)&d_A, sizeof(double) * Width * Width);
    cudaMalloc((void**)&d_B, sizeof(double) * Width);
    cudaMalloc((void**)&d_X, sizeof(double) * Width);

    cudaMemcpy(d_A, A, sizeof(double) * Width * Width, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, sizeof(double) * Width, cudaMemcpyHostToDevice);  


    do_parallel_matmul<<<dimB, dimT>>>(d_A, d_B, d_X, Width);   
    

    cudaMemcpy(X, d_X, sizeof(double) * Width, cudaMemcpyDeviceToHost);

    finish = clock();
    
    totaltime=(double)(finish-start)/CLOCKS_PER_SEC;   

    printf("%f", totaltime);

这个时间比连续时间长得多，如下所示：

clock_t start,finish;
double totaltime;
start = clock(); 

do_seq_matmult();

finish = clock();
    
totaltime=(double)(finish-start)/CLOCKS_PER_SEC;   

printf("%f", totaltime);

所以我不知道我是否应该只测量CUDA内核时间如下：

clock_t start,finish;
double totaltime;
start = clock(); 

do_parallel_matmul();

finish = clock();
    
totaltime=(double)(finish-start)/CLOCKS_PER_SEC;   

printf("%f", totaltime);

并避免主机和设备之间的内存复制...

我之所以问上述问题是因为我必须提交并行执行和顺序执行之间的比较...但是如果我在 CUDA 中测量内存副本，那么 CUDA 和 C 之间并没有很好的区别...

编辑：

void do_seq_matmult(const double *A, const double *X, double *resul, const int tam)
{
    *resul = 0;
    for(int i = 0; i < tam; i++)
    {
        for(int  j = 0; j < tam; j++)
        {
            if(i != j)
                *resul += A[i * tam + j] * X[j];
        }
    }
}

__global__ void do_parallel_matmul( double * mat_A, 
                            double * vec, 
                            double * rst, 
                            int dim)
{
     int rowIdx = threadIdx.x + blockIdx.x * blockDim.x; // Get the row Index 
     int aIdx;
     while( rowIdx < dim)
     {
          rst[rowIdx] = 0; // clean the value at first
          for (int i = 0; i < dim; i++)
          {
               aIdx = rowIdx * dim + i; // Get the index for the element a_{rowIdx, i}
               rst[rowIdx] += (mat_A[aIdx] * vec[i] ); // do the multiplication
          }
          rowIdx += gridDim.x * blockDim.x;
     }
     __syncthreads();
}

Answer 1

您在测量时使用了错误的函数。 clock 测量您的进程在 CPU 上花费的时间，而不是挂钟时间。

Answer 2

看看High Precision Timer lib，它使用操作系统相关的计时函数来测量时间。

它使用一组函数，可以为您提供微秒精度。

如果你在 Windows 上，你应该使用 QueryPerformanceFrequency 和 QueryPerformanceCounter 在 Linux 上：gettimeofday()

它非常轻巧且易于使用。适用于 windows 和 linux。

Answer 3

一些想法：

1. 对设备内存的分配进行计时并与没有主机分配内存的 CPU 进行比较是不公平的。

2. 如果 cudaMalloc((void**)&d_A, sizeof(double) * Width * Width); 是第一个 CUDA 调用，它将包括 CUDA 上下文创建，这可能是一个很大的开销。

3. 计时 cudamemcpy 不是一个公平的 CPU/GPU 比较，因为这个时间将取决于系统的 PCI-e 带宽。另一方面，如果您从 CPU 的角度将内核视为加速，则需要包含 memcpy。为了使 PCI-e 带宽达到峰值，请使用页面锁定内存。

4. 如果您的应用程序要多次运行乘法运算，那么您可以通过将副本与内核执行重叠来隐藏大部分 memcpy。这在配备双 DMA 引擎的 Tesla 设备上效果更好。

5. 为内核本身计时需要您在停止计时器之前将 CPU 与 GPU 同步，否则您将只为内核启动本身而不是执行计时。从 CPU 调用内核是异步的。如果您想为 GPU 上的内核执行计时，请使用 cudaEvents。

6. 在 GPU 上运行多个线程以获得公平的比较。

7. 改进内核，你可以做得更好。