cuda - GPU 中每个操作的时钟周期数

Question

有什么方法可以找到使用 CUDA 在 GPU 中执行不同操作(如除法、减法和加法)所需的时钟周期数？

Answer 1

您可能想看看时钟设备功能:clock .
当延迟和吞吐量开始发挥作用时，它可能不是您提到的每条指令周期的基本事实指标，但绝对是一个有用的工具。

这是一个关于如何将其用于吞吐量估计的代码示例:

__global__ void timing_1(int N, float4* input, float4* output)
{
    float4 a,b,c,d ;
    a = input[0]; b = input[1]; c = input[2]; d = input[3];

    long long int start = clock64();

    for (int k = 0 ; k < N ; ++k)
    {
        a.x = 1.0 / a.x ; a.y = 1.0 / a.y ; a.z = 1.0 / a.z ; a.w = 1.0 / a.w ;
        b.x = 1.0 / b.x ; b.y = 1.0 / b.y ; b.z = 1.0 / b.z ; b.w = 1.0 / b.w ;
        c.x = 1.0 / c.x ; c.y = 1.0 / c.y ; c.z = 1.0 / c.z ; c.w = 1.0 / c.w ;
        d.x = 1.0 / d.x ; d.y = 1.0 / d.y ; d.z = 1.0 / d.z ; d.w = 1.0 / d.w ;
    }

    long long int stop = clock64();

    // make use of data so that compiler does not optimize it out
    a.x += b.x + c.x + d.x ;
    a.y += b.y + c.y + d.y ;
    a.z += b.z + c.z + d.z ;
    a.w += b.w + c.w + d.w ;

    output[threadIdx.x + blockDim.x * blockIdx.x] = a ;

    if (threadIdx.x == 0)
        ::printf ("timing_1 - Block [%d] - cycles count = %lf - cycles per div = %lf\n", blockIdx.x, ((double)(stop - start)), ((double)(stop-start))/(16.0*(double)N)) ;
}

对于延迟，您希望在计算之间具有依赖性:

__global__ void timing_2(int N, float4* input, float4* output)
{
    float4 a  ;
    a = input[0]; 

    long long int start = clock64();

    for (int k = 0 ; k < N ; ++k)
    {
        a.y = 1.0 / a.x ; a.z = 1.0 / a.y ; a.w = 1.0 / a.z ; a.x = 1.0 / a.w ;
    }

    long long int stop = clock64();

    output[threadIdx.x + blockDim.x * blockIdx.x] = a ;

    if (threadIdx.x == 0)
        ::printf ("timing_2 - Block [%d] - cycles count = %lf - cycles per div = %lf\n", blockIdx.x, ((double)(stop - start)), ((double)(stop-start))/(4.0*(double)N)) ;
}

您希望使用每个 SM 的少量线程和块运行它，以避免计算重叠，这会使您的挂钟计时器与单个计算不一致。

对于 GTX 850m 上的 32 个线程和 5 个块，我获得了每除法 128 个周期的吞吐量和 142 个周期的延迟，使用常规数学以单精度计算 - 转换为函数调用(nvcc 7.5，sm_50)。

使用快速数学时，我得到 2.5 个周期的吞吐量和 3 个周期的延迟。