AnsweredAssumed Answered

memory conflicts when reading same address

Question asked by firespot on Aug 9, 2014
Latest reply on Aug 20, 2014 by dipak



The AMD SDK 2.9 states ( "Under certain conditions, one unexpected case of a channel conflict is that reading from the same address is a conflict, even on the FastPath. This does not happen on the read-only memories, such as constant buffers, textures, or shader resource view (SRV); but it is possible on the read/write UAV memory or OpenCL global memory. ..."


What precisely counts as read-only memory / constant buffer?


Suppose there are 2 kernel variants:


__kernel void kernelA(__constant int const * someMem);

__kernel void kernelB(__global int const * someMem);


Does someMem qualify in both cases as read-only with respect to because the type declaration is always int const *, or is it here per se only the former case as the memory is  __constant qualified?


If the latter kernel buffer does not yet count already as read-only given the kernel-declaration, will it count then as read-only if the underlying buffer has been constructed with the CL_MEM_READ_ONLY flag?


The question is for the following reason:

One kernel calculates values in the buffer (write-only), while several other kernels subsequently only make read-only access to that buffer, where it frequently happens that the same address is accessed (no way to avoid that algorithmically speaking). I cannot declare the buffer itself CL_MEM_READ_ONLY. If kernelB above would, based on its declaration itself, not already count as read-only would the following pattern avoid same-address memory conflicts:

-) create two buffers, one with CL_MEM_WRITE_ONLY flag, one with CL_MEM_READ_ONLY

-) the kernel making the write-only operations writes its values into the former buffer

-) copy the write-only buffer to the read-only buffer

-) all other kernels access read-only buffer


Last question: Suppose the raw data is simply an array of ints, but in kernel-code are fetched using vector types (e.g. int16), does it count as "same address" if any scalar int memory is shared between any two vector reads, or does only the start address (the first of all 16 scalar ints) matter?