Intel® Quartus® Prime Software
Intel® Quartus® Prime Design Software, Design Entry, Synthesis, Simulation, Verification, Timing Analysis, System Design (Platform Designer, formerly Qsys)
All support for Intel NUC 7 - 13 systems has transitioned to ASUS. Read latest update.
16502 Discussions

OpenCL burst Non-aligned access.


I have a OpenCL code like bellow,

the report.html says that I have burst non-aligned access to Z,

but I have read 4x16xchar = 512 bit every clock,

why this is not aligned?

typedef struct{ char ww[16]; }packed_16;   typedef struct{ packed_16 w[4]; }packed_4_16;   __kernel void Read( __global const packed_16* restrict Z, int end, ){ for(int s=0 ; s<end; s++){   // some statement about a_address, b_address // packed_4_16 data; #pragma unroll for(int i=0; i<4; i++){ packed_16 Z_cache = Z[ a_address + b_address + i ]; data.w[i] = Z_cache; } } }

0 Kudos
1 Reply
Valued Contributor III

Personally, I have never seen coalesced accesses be inferred as aligned ports, they have always been non-aligned in my experience. However, this might not be universal case and for very regular coalesced accesses, the compiler might infer aligned ports. In your case, you have two offsets in your address (a_address and b_address); it is likely that it is very difficult for the compiler to predict whether these offsets would be aligned or not and hence, it infers non-aligned ports. In practice, you are not going to lose much performance, if at all, if all your accesses are aligned but the memory port is non-aligned.

0 Kudos