OpenCL burst Non-aligned access.

HHo2 · ‎12-27-2018

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;
		}
	}
}

HRZ · ‎12-27-2018

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.