Is my GTX680 really performing

4
votes

I am trying to test my GTX680 for calculation performance as I have some doubts how much it is really performing. I was wondering if maybe somebody can also test on his GTX 680 if same results are given or maybe tell me what it can be done better as to tap more performance from the card

I've written this small program

#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);
}

Compiling with nvcc I get this 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; 
 }

The kernel runs in 1.936ms

My calculations show that GFLOPS performance was 1.1 TFLOP just a third of the theoretical value of 3TFLOPS (Ref: http://www.geforce.com/hardware/desktop-gpus/geforce-gtx-680 ).. Why is it so slow?

Details of my calculations are as follows

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)

Total FLOPS for 1 iteration in the loop is 7 FLOPS

Considering only iterations

We have 300 iterations per thread We have 1024*1000 Blocks

Total iterations FLOPS = 300*1024*1000*7 = 2.15 GFLOPS

Total kernel time is 1.936ms

Hence throughput = 1.11 TFLOPS

Thanks in advance for help

Daniel

4
cuda
Can't help you, this is out of my reach, but it is damn impressive.. - Simon Carlson
Trying unrolling the loop lots, the tight loop is consuming a lot of cycles as it is. - talonmies
Do you mean that conditional jumping ( @%p2 bra ) takes more the 2 flop time to complete...? I am doing some other research..Is integer arithmetic slower in gpus as I am using gpu-z and it's telling me that it can only reach 500giops/s while it reaches nearly 2tflops? - Daniel
In your question you are calculating the number of instructions executed. These should be counted from the SASS (assembly code) not the intermediate PTX code. Use cuobjdump -sass EXECUTABLE/SO to get the SASS code. The CUDA command line profiler, nvprof, Visual Profiler, and Nsight VSE CUDA Profiler can all collect inst_executed counter. Nsight VSE CUDA Profiler provides the experiment "Achieved Flops" if you are interested in counting just floating point operations (not instructions). - Greg Smith

4 Answers

3
votes

This example program builds on @Robert Crovella answer. Robert's kernel is limited by data dependencies. By reducing data dependencies between FMA instructions this kernel should achieve 2.4-2.5 TFLOPS on a GTX680.

The current implementation is instruction fetch and data dependencies limited. The kernel should be able to be tweaked to improve Achieved FLOPS by another 10%.

Nsight Visual Studio Edition 2.x and the new 3.0 RC candidate provide the metrics you need to analysis this kernel.

In 2.x and 3.0 you should use the following experiments to analyze the kernel:

  1. Instruction Statistics - SM Activity - Verify that all SMs are close to 100%
  2. Issue Efficiency - Eligible Warps - On Kepler Eligible Warps Per Active Cycle has to be greater than 4 in order for each warp scheduler to issue an instruction every cycle.
  3. Issue Efficiency - Issue Stalls - Warp Issue Efficiency will specify how often each warp scheduler was unable to issue due to insufficient number of eligible warps. If this is high then the Issue Stall Reasons will help identify the limiter.
  4. Achieved FLOPs - This displays both a breakdown of the type and rate of single and double precision floating point operations performed by the kernel.

In the case of Robert's kernel Execution Dependencies was extremely high as each instruction had a read after write dependency. By increasing the instruction level parallelism we tripled the performance. The kernel is now primarily instruction fetch limited.

The new Nsight VSE 3.0 RC (available today) will also show the assembly or source code annotated with per instruction statistics such as the number of instructions executed and the number of active threads per instruction. In this example the tool can be used to identify the data dependencies and make sure that the compiler is generating FMA instructions which are required to reach greater than 50% theoretical achieved 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);
}
1
votes

I think the problem is that you're using an integer multiply. 32-bit integer multiply on the compute capability 3.0 architecture is only 1/6th of the 32-bit floating point throughput (refer the table below, taken from the CUDA C Programming Guide version 5.5). Compare the 32-bit integer multiply performance with the 32-bit floating point performance for the 3.0 architecture.

Some other integer operations and type conversions that are mostly used in compute applications have similarly reduced performance on 3.0.

enter image description here

0
votes

See if you get better results with this code. It's just an example, it doesn't do precisely the same thing as your code, and I think you'll have to recount flops.

#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);
}

When I compile this with arch = sm_20 or sm_30, I get 10 fma instructions in a row in the kernel loop, with no intervening code. I think it will run faster, and closer to peak theoretical flops than your code. Yes, there is a difference between integer OPs/second and floating point OPs/second. If you do run this code, please comment back and let me know what your calculated perf was.

0
votes

Your test kernel is doing integer operations, rather than floating point operations. So FLOPS is the wrong metric for that kernel all together.

FLOPS = FLoating point Operations Per Second

Back to the original question though, your kernel is slow because GPUs are optimized for floating point calculations, rather than integer calculations.

To do a proper test, try converting your test kernel to use floating point numbers, rather than integers.

Additionally, in the loop where you are annotating FLOPS to the steps, FLOPS again doesn't make sense, since it is a per second measure and those are integer operations. Once you've converted it, just count them as individual floating point operations, not floating point operations per second.