Intel® Quartus® Prime Software
Intel® Quartus® Prime Design Software, Design Entry, Synthesis, Simulation, Verification, Timing Analysis, System Design (Platform Designer, formerly Qsys)
Announcements
Intel Support hours are Monday-Fridays, 8am-5pm PST, except Holidays. Thanks to our community members who provide support during our down time or before we get to your questions. We appreciate you!

Need Forum Guidance? Click here
Search our FPGA Knowledge Articles here.
14944 Discussions

OpenCL burst Non-aligned access.

HHo2
Beginner
993 Views

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
HRZ
Valued Contributor II
88 Views

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.

Reply