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 个块，通过常规数学转换为函数调用（nvcc 7.5，sm_50），我得到了每个分区 128 个周期的吞吐量和 142 个周期的单精度延迟。

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