The Cuda 2.0 manual (p. 62) says
“A common conflict-free case is when all threads of a half-warp read from an address within the same 32-bit word.”
However, I would like to guarantee that if the same location in memory is read by two different threads, then the reads will be serialized.
Is there a way to ensure this?
TIA
Out of curiosity, why would you want this? Since reads don’t change the data, the ordering of them is irrelevant. Are you wanting to simulate bank conflicts on purpose?
It is an application where data can be read and written to the same location in shared memory. Just want to ensure that it will work correctly. The writes will serialize and ensure correctness, and want to be sure that what is being read is deterministic, by having that serialized
Im not sure i understand all the details of what you want to do, but this seems pretty dangerous as you cannot make any assumptions on the order in which the threads will be launched.
This is precisely what __syncthreads() is for. If you put it between the write and the read operation, you will ensure that the read operation will wait until the previous write is complete.