AnsweredAssumed Answered

Difference in accessing global memory with (int2*) and (int*)?

Question asked by viscocoa on Mar 22, 2012
Latest reply on May 19, 2012 by viscocoa

I just used a whole day to fix a small problem. My kernel is conceptually like below:

 

__kernel

void myKernel(__global int2* data)

{

     do something, and get two int values a and b

     data[a].x = something;

     data[b].y = something_else;

}

 

The above code works well on CPUs. On HD5870, in very rare and random cases, one of the two writes to global memory does not work. The data to be written is lost, and the value in data[] is not changed.

I then changed the code to:

 

__kernel

void myKernel(__global int* data)

{

     do something, and get two int values a and b

     data[a<<1] = something;

     data[(a<<1)+1] = something_else;

}

 

The above code work well on HD5870!

 

-------------------------------------------------------------------------

It looks like writing to the two components of an int2 simultaneously (maybe from two compute units) causes a conflict, and one of them will overwrite the other. However, writing to int is independent.

 

I wonder if this is a bug of the OpenCL implementation.

 

The above logic is used in many circumstances like radix sort. Hope it is helpful for those who has the same problem.

Outcomes