Opencl, Will using global pointer to different local memory create branching?

user_4863746 0
Viewed 29

I have a kernel like this

kernel void k1(global int * a, global int * b, global int * c){
  int i = get_local_id(0);
  global int* ptr;
  if (i==0) ptr = &a[0];
  if (i==1) ptr = &b[0];
  if (i==2) ptr = &c[0];
  if (i>2) return;
  ptr[0]++;

}

My question is, will the line ptr[0]++ create branching in accessing memory? or will this save time compared to a[0]++;b[0]++;c[0]++; within same workitem.

Thanks for your time.

opencl kernel branch optimization
1 Answers

No, this does not create branching, but it will break coalesced memory access at least in some workgroups.

You can also entirely replace the if(i==0) ptr = &a[0]; etc. by either ternary operators (condition ? true_action : false_action) or by bit masking. Such simple if statements and the ternary operator are most likely to be converterd to branchless assembly. If it weren't for the guard clause (if(i>2) return;), the entire code then is branchless.

Although there is no branching, more importantly, especially with more complex pointer arithmetic than you have (for example when a lookup table determines the pointer value), the arithmetic telling which memory location should be accessed will break coalesced memory access in some workgroups, which could lead to much more slowdown compared to simple branching itself (although branching wouldn't resolve the misaligned memory access as well).

In some applications it is even possible that loading more global memory in coalesced manner is faster than reducing global memory access but breaking coalescence in the process.

user_91789920
Related