4

我正在尝试测试我的 GTX680 的计算性能，因为我怀疑它的实际性能如何。我想知道如果给出相同的结果，是否有人也可以在他的 GTX 680 上进行测试，或者告诉我可以做得更好以从卡中获得更多性能

我写了这个小程序

#include <stdlib.h>
#include <stdio.h>
#include <float.h>
#include "cuComplex.h"
#include "time.h"
#include "cuda_runtime.h"
#include <iostream>
using namespace std;
__global__ void test(int loop, int *out)
{
    register int a=0;
    for (int x=0;x<loop;x++)
    {
        a+=x*loop;
    }


    if (out!=NULL) *out=a;


}
int main(int argc, char *argv[])
{
    float timestamp;
    cudaEvent_t event_start,event_stop;
    // Initialise
    cudaDeviceReset();
    cudaDeviceReset();
    cudaSetDevice(0);
    cudaThreadSetCacheConfig(cudaFuncCachePreferShared);        
    // Allocate and generate buffers
    cudaEventCreate(&event_start);
    cudaEventCreate(&event_stop);
    cudaEventRecord(event_start, 0);
    dim3 threadsPerBlock;
    dim3 blocks;
    int b=1000; 
    threadsPerBlock.x=32;
    threadsPerBlock.y=32;
    threadsPerBlock.z=1;
    blocks.x=1;
    blocks.y=1000;
    blocks.z=1;

    test<<<blocks,threadsPerBlock,0>>>(300,
            NULL
            );

    cudaEventRecord(event_stop, 0);
    cudaEventSynchronize(event_stop);
    cudaEventElapsedTime(&timestamp, event_start, event_stop);
    printf("Calculated in %f", timestamp);
}

用 nvcc 编译我得到这个 PTX

//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sat Sep 22 02:35:14 2012 (1348274114)
// Cuda compilation tools, release 5.0, V0.2.1221
//

.version 3.1
.target sm_30
.address_size 64

    .file   1 "/tmp/tmpxft_00000e7b_00000000-9_perf.cpp3.i"
    .file   2 "/opt/home/daniel/a/perf.cu"

 .visible .entry _Z4testiPi(
    .param .u32 _Z4testiPi_param_0,
    .param .u64 _Z4testiPi_param_1
 )
 {
    .reg .pred      %p<4>;
    .reg .s32       %r<15>;
    .reg .s64       %rd<3>;


    ld.param.u32    %r6, [_Z4testiPi_param_0];
    ld.param.u64    %rd2, [_Z4testiPi_param_1];
    cvta.to.global.u64      %rd1, %rd2;
    mov.u32         %r13, 0;
    .loc 2 12 1
    setp.lt.s32     %p1, %r6, 1;
    mov.u32         %r14, %r13;
    mov.u32         %r11, %r13;
    @%p1 bra        BB0_2;

 BB0_1:
    .loc 2 14 1
    mad.lo.s32      %r14, %r11, %r6, %r14;
    .loc 2 12 20
    add.s32         %r11, %r11, 1;
    .loc 2 12 1
    setp.lt.s32     %p2, %r11, %r6;
    mov.u32         %r13, %r14;
    @%p2 bra        BB0_1;

 BB0_2:
    .loc 2 18 1
    setp.eq.s64     %p3, %rd2, 0;
    @%p3 bra        BB0_4;

    .loc 2 18 1
    st.global.u32   [%rd1], %r13;

 BB0_4:
    .loc 2 21 2
    ret; 
 }

内核运行时间为 1.936ms

我的计算表明 GFLOPS 性能为 1.1 TFLOP，仅为 3TFLOPS 理论值的三分之一（参考：http ://www.geforce.com/hardware/desktop-gpus/geforce-gtx-680 ）.. 为什么这么慢?

我的计算细节如下

mad.lo.s32      %r14, %r11, %r6, %r14;  //2 FLOPS
.loc 2 12 20
 add.s32         %r11, %r11, 1;     //1 FLOP
.loc 2 12 1
 setp.lt.s32     %p2, %r11, %r6;    //1 FLOP
 mov.u32         %r13, %r14;        // 1 FLOP
 @%p2 bra        BB0_1;             //1 FLOP

 + 1 FLOP (just as a buffer as I don't know branching how much it takes)

循环中 1 次迭代的总 FLOPS 为 7 FLOPS

只考虑迭代

我们每个线程有 300 次迭代 我们有 1024*1000 个块

总迭代 FLOPS = 300*1024*1000*7 = 2.15 GFLOPS

总内核时间为 1.936ms

因此吞吐量 = 1.11 TFLOPS

提前感谢您的帮助

丹尼尔

4

4 回答 4

3

这个示例程序建立在@Robert Crovella 的答案之上。Robert 的内核受到数据依赖性的限制。通过减少 FMA 指令之间的数据依赖性，这个内核应该在 GTX680 上实现 2.4-2.5 TFLOPS。

当前的实现是指令获取和数据依赖性受限。内核应该能够被调整以将实现的 FLOPS 再提高 10%。

Nsight Visual Studio Edition 2.x 和新的 3.0 RC 候选版本提供了分析此内核所需的指标。

在 2.x 和 3.0 中，您应该使用以下实验来分析内核：

  1. 指令统计 - SM 活动 - 验证所有 SM 是否接近 100%
  2. 发布效率 - 合格 Warps - 在 Kepler 上 每个活动周期的合格 Warps 必须大于 4，以便每个 warp 调度程序在每个周期发出一条指令。
  3. 问题效率 - 问题停顿 - Warp 问题效率将指定每个 warp 调度程序由于合格的 warp 数量不足而无法发出的频率。如果这个值很高，那么问题失速原因将有助于识别限制器。
  4. 实现的 FLOPs - 这显示了内核执行的单精度和双精度浮点运算的类型和速率的细分。

在 Robert 的内核中，执行依赖性非常高，因为每条指令都具有读写依赖性。通过增加指令级并行性，我们将性能提高了两倍。内核现在主要是指令获取受限。

新的 Nsight VSE 3.0 RC（今天可用）还将显示带有每条指令统计注释的程序集或源代码，例如执行的指令数和每条指令的活动线程数。在此示例中，该工具可用于识别数据依赖性并确保编译器正在生成 FMA 指令，这些指令需要达到大于 50% 的理论实现 FLOPS。

__global__ void test(float loop, float *out)
{
    register float a=1.0f;
    register float b=1.0f;
    register float c=1.0f;
    register float d=1.0f;
    register float e=1.0f;
    register float f=1.0f;
    register float g=1.0f;
    register float h=1.0f;

    for (float x=0;x<loop;x++)
    {
        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;

        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;

        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;

        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;

        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;
    }
    if (out!=NULL) *out=a+b+c+d+e+f+g+h;
}

int main(int argc, char *argv[])
{
    float timestamp;
    cudaEvent_t event_start,event_stop;
    // Initialise
    cudaDeviceReset();
    cudaSetDevice(0);
    cudaThreadSetCacheConfig(cudaFuncCachePreferShared);
    // Allocate and generate buffers
    cudaEventCreate(&event_start);
    cudaEventCreate(&event_stop);
    cudaEventRecord(event_start, 0);
    dim3 threadsPerBlock;
    dim3 blocks;
    threadsPerBlock.x=32;
    threadsPerBlock.y=32;
    threadsPerBlock.z=1;
    blocks.x=1;
    blocks.y=1000;
    blocks.z=1;

    test<<<blocks,threadsPerBlock,0>>>(30,NULL);

    cudaEventRecord(event_stop, 0);
    cudaEventSynchronize(event_stop);
    cudaEventElapsedTime(&timestamp, event_start, event_stop);
    printf("Calculated in %f\n", timestamp);
}
于 2012-10-30T03:10:14.623 回答
1

我认为问题在于您使用的是整数乘法。计算能力 3.0 架构上的 32 位整数乘法仅是 32 位浮点吞吐量的 1/6（请参阅下表，取自 CUDA C 编程指南 5.5 版）。将 32 位整数乘法性能与 3.0 架构的 32 位浮点性能进行比较。

其他一些主要用于计算应用程序的整数运算和类型转换在 3.0 上也同样降低了性能。

在此处输入图像描述

于 2013-11-13T06:51:06.357 回答
0

您的测试内核正在执行整数运算，而不是浮点运算。所以 FLOPS 是该内核的错误指标。

FLOPS = FLoating point Operations Per Second

回到最初的问题，您的内核很慢，因为 GPU 针对浮点计算进行了优化，而不是整数计算。

要进行正确的测试，请尝试将测试内核转换为使用浮点数，而不是整数。

此外，在将 FLOPS 注释到步骤的循环中，FLOPS 再次没有意义，因为它是每秒测量值，并且是整数运算。转换后，只需将它们计为单独的浮点运算，而不是每秒的浮点运算。

于 2012-10-29T16:23:44.747 回答
0

看看您是否使用此代码获得了更好的结果。这只是一个例子，它与你的代码做的事情并不完全相同，我认为你必须重新计算失败。

#include <stdio.h>
using namespace std;
__global__ void test(float loop, float *out)
{
    register float a=1.0f;
    for (float x=0;x<loop;x++)
    {
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
    }


    if (out!=NULL) *out=a;


}
int main(int argc, char *argv[])
{
    float timestamp;
    cudaEvent_t event_start,event_stop;
    // Initialise
    cudaDeviceReset();
    cudaSetDevice(0);
    cudaThreadSetCacheConfig(cudaFuncCachePreferShared);
    // Allocate and generate buffers
    cudaEventCreate(&event_start);
    cudaEventCreate(&event_stop);
    cudaEventRecord(event_start, 0);
    dim3 threadsPerBlock;
    dim3 blocks;
    threadsPerBlock.x=32;
    threadsPerBlock.y=32;
    threadsPerBlock.z=1;
    blocks.x=1;
    blocks.y=1000;
    blocks.z=1;

    test<<<blocks,threadsPerBlock,0>>>(30,
            NULL
            );

    cudaEventRecord(event_stop, 0);
    cudaEventSynchronize(event_stop);
    cudaEventElapsedTime(&timestamp, event_start, event_stop);
    printf("Calculated in %f\n", timestamp);
}

当我用 arch = sm_20 或 sm_30 编译它时，我在内核循环中连续获得 10 条 fma 指令，没有中间代码。我认为它会比你的代码运行得更快，更接近理论失败的峰值。是的，整数 OPs/秒和浮点 OPs/秒之间存在差异。如果您确实运行此代码，请回复评论并让我知道您计算的性能是什么。

于 2012-10-29T16:11:50.513 回答