Huge difference in memory bandwidth for IVB

Biao_W_ — Fri, 03 Jan 2014 23:32:29 GMT

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

Hi, all:

Biao_W_ — Tue, 07 Jan 2014 22:29:22 GMT

Hi, all:

Answers are highly appreciated.

Hi Biao,

Raghupathi_M_Intel — Thu, 09 Jan 2014 23:16:37 GMT

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

Quote:Raghu Muthyalampalli

Biao_W_ — Thu, 09 Jan 2014 23:36:36 GMT

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

topic Hi, all: in OpenCL* for CPU

Huge difference in memory bandwidth for IVB

Hi, all:

Hi Biao,

Quote:Raghu Muthyalampalli