Higher Tested Bandwidth with Caches Compared to Theoretical Memory Bandwidth in GPU when Accessing Global Memory

I am currently working on optimizing memory access in a GPU-accelerated application and have encountered an intriguing scenario. When I use linear read in kernel code while accessing global memory, I have noticed that the tested memory bandwidth appears to be higher than the theoretical memory bandwidth. I'm seeking to understand the possible reasons behind this behavior.

Here is my kernel code

__kernel void read_single(__global DATATYPE *input,__global DATATYPE *output)
{
        DATATYPE val = (DATATYPE)(0.0f);
        IDXTYPE gid = get_global_id(0);

        val = val + input[0];
        val = val + input[1];
        val = val + input[2];
        val = val + input[3];
        val = val + input[4];
        val = val + input[5];
        val = val + input[6];
        val = val + input[7];
        val = val + input[8];
        val = val + input[9];
        val = val + input[10];
        val = val + input[11];
        val = val + input[12];
        val = val + input[13];
        val = val + input[14];
        val = val + input[15];
        val = val + input[16];
        val = val + input[17];
        val = val + input[18];
        val = val + input[19];
        val = val + input[20];
        val = val + input[21];
        val = val + input[22];
        val = val + input[23];
        val = val + input[24];
        val = val + input[25];
        val = val + input[26];
        val = val + input[27];
        val = val + input[28];
        val = val + input[29];
        val = val + input[30];
        val = val + input[31];

        output[gid] = val;
}

I used GPU Navi 14 [Radeon Pro W5500]

• Memory Size: 8 GB
• Memory Type: GDDR6
• Memory Bus : 128 bit
• Bandwidth : 224.0 GB/s

But my test results is

• Global Memory Read: Single
• Size (Bytes) 33554432
• Avg. Kernel Time (sec) 1.56687e-05
• Avg Bandwidth (GBPS) 2141.5

I did global memory bandwidth test and expected that it will be lower than theoretical value

OpenCL is something of a black box to OpenCL programmers. Who knows what it actually does inside the box. I have some thoughts on your test though.

Like when you compile a c program, the compiler will look for ways to optimise the computation. Sometimes, this will involve re-ordering code or evaluating the return value of functions if the input is known at compile time. One obvious optimisation that could be made with your test is from observing that most of the kernel calculation (the 32 val+=... statements) is the same in each thread, so only has to be evaluated at runtime once, not once per thread. It certainly doesn't have to read more than the first 32 elements of the input array. Then all it has to do is copy that value 'global-size' times to the output. I think if you inspect your output buffer, every element will hold the same value.

You didn't say what your global-size was, but let's say it was 250,000. Then, your output buffer would be about 1MB (assuming your DATATYPE is 32 bit float). You would then have updated 1MB in 1.57e-5 seconds, which gives a bandwidth of approx 63.7 GB/s.

If you have 32MB for your input and want to read all the way through, you need to be using different indices on each thread. You could do something like

private DATATYPE *threadindexstart = input + 32*gid;
val = *threadindexstart;
val += *(threadindexstart + 1);...