Huge difference in memory bandwidth for IVB

[Biao_W_]

Hi, 

I used memory bandwidth benchmark in AMD APP 2.9 SDK and test the memory bandwidth in IVB HD 4000, the original benchmark only test the data type of "float" while I extend it to int, short and unsigned char, all benchmark using scalar data type, no vector type is used. I observed a huge difference of memory bandwidth between 32-bit memory access and none 32-bit memory access, following are my results:

Global Memory Read: Single (all threads read from a single memory location) GB/s in Ivy bridge HD 4000

Float     Int      Short   uchar 

8.79     8.79     0.60     0.60

Comparatively, I run the same benchmark on Haswell i4770S HD4600 graphics and the bandwidth gap shrinks significantly.

Float         Int        Short    uchar 

56.05       57.60     33.90    31.91

The testing kernel code is presented as follows, where the DATATYPE can be defined during the online compilation:

__kernel void read_single(__global DATATYPE *input,__global DATATYPE *output)
{
    DATATYPE val = (DATATYPE)(0.0f);
    uint 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;
}

The question are,

1. why are there so much memory bandwidth difference between 32 bit and none 32bit data type in  IVB?

2. why haswell doing much better when access none 32 bit data types?

Best

[Biao_W_]

Hi, all:

Answers are highly appreciated.

[Raghupathi_M_Intel]

Hi Biao,

Haswell has several memory access performance improvements for these data types. That's why you see the huge BW performance difference going from IVB to HSW. So what you are seeing is expected.

Raghu

[Biao_W_]

Raghu Muthyalampalli (Intel) wrote:

Hi Biao,

Haswell has several memory access performance improvements for these data types. That's why you see the huge BW performance difference going from IVB to HSW. So what you are seeing is expected.

Raghu

Hi, Raghu:

Thanks for your reply.

However, could you provide more detail about the data type improvements on HSW? Or give me some links to related documents? 

Thanks