Stream benchmark

The stream benchmark aims at measuring memory bandwidth performance of the architecture. We simply add two vectors, in what is our “Hello World” program. Simplified version of the C# code can be found on our SDK on github.

As for Expm1 benchmark we run the generated code against the best handwritten native code we could write. This allows us to compare performance differences between hybridized code and handwritten native code using intrinsics.

public static void Run(int N, double[] a, double[] b)
    int start = threadIdx.x + blockDim.x * blockIdx.x;
    int step = blockDim.x * gridDim.x;
    for(int k = start; k < N; k += step) 
        a[k] += b[k];

No Hint

In a first version, we don’t give any hint about alignment or index coalescence. This has little to no effect on CUDA machines, and hardware prefetch looks doing a very good job on Intel Xeon machines. However, this leads to a gather instruction, which has a little performance penalty associated on Xeon Phi. We’ll see later how to restore performance.

Architecture Generated Handwritten Ratio
NVIDIA- P100 479.9 495.4 96.8%
NVIDIA – K20C 167 185 90.2%
INTEL – XEON PHI – 7210 327.7 381.7 85.8%
INTEL – Xeon E5 1620 v3 – 3.5 GHz 34.4 35 98.3%

Coalesced – Aligned

In a second version, we specify that our backend storage is an aligned memory location, and that index will stay coalesced and aligned:

public static void StreamDouble(alignedstorage_double a, alignedstorage_double b, int offset, int count)
    alignedindex start = offset + threadIdx.x + blockIdx.x * blockDim.x ;
    alignedindex end = count + offset;
    for (alignedindex i = start; i < end; i += blockDim.x * gridDim.x)
        a[i] = a[i] + b[i];

Architecture Generated Handwritten Ratio
INTEL – XEON PHI – 7210 373.8 381.7 97.9%

Tags: , ,