Forum Discussion

HHo2's avatar
HHo2
Icon for New Contributor rankNew Contributor
7 years ago

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

1 Reply

  • HRZ's avatar
    HRZ
    Icon for Frequent Contributor rankFrequent Contributor
    7 years ago

    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.