OpenCL* for CPU
Ask questions and share information on Intel® SDK for OpenCL™ Applications and OpenCL™ implementations for Intel® CPU
Announcements
This forum covers OpenCL* for CPU only. OpenCL* for GPU questions can be asked in the GPU Compute Software forum. Intel® FPGA SDK for OpenCL™ questions can be ask in the FPGA Intel® High Level Design forum.
1663 Discussions

Huge difference in memory bandwidth for IVB

Biao_W_
Beginner
136 Views

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

0 Kudos
3 Replies
Biao_W_
Beginner
136 Views

Hi, all:

Answers are highly appreciated.

 

Raghupathi_M_Intel
136 Views

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_
Beginner
136 Views

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

Reply