Archives Discussions

[firespot]

Hi,

The AMD SDK 2.9 states (6.1.1.2): "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 6.1.1.2 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?

Thanks!

[dipak]

Hi,

__constant and const are totally different things. In case of __constant, the data is in constant memory which may be physically separated from global memory on device and accessing one type of memory using another address qualifier is illegal. Where as "const" is a compiler level thing, and you can actually cast it away.

For example, you create a buffer with read-only mode and pass it to a kernel having argument type as constant for that argument. In this case, data will be loaded into constant memory transparently by the run-time and the accessing will be faster. You can also directly access the read only buffer via "global const *" from another kernel but the access path will be different.

You can cast pointer of a scalar data type to a pointer of vector data type as long as they agree with the data type alignment rule (for details see

Chapter 6.2 Conversions and Type Casting in The OpenCL Specification1.2)

[firespot]

Hi, well I am aware of the hardware-related issues between __global and __constant. Indeed that hardware difference is part of my problem because the data for which same-address reads occur may by no means fit into __constant memory (so practically speaking that's out of scope); so that's why my kernel-parameter is __global qualified, yet I can declare the elements themselves to be const and can make the buffer (or a copy-instance of it, see above) CL_MEM_READ_ONLY.

The prime question however yet remains: Waht counts as read-only in 6.1.1.2? I can come up with three reasonable candidates:

1) memory is __constant qualified

2) elements memory points to are const-qualified

3) buffer is declared CL_MEM_READ_ONLY.

[any else?]

Does only the first of them avoid channel conflicts as different hardware resources are used, or is it also any of the latter two?

thanks !

[maxdz8]

I'd add

4) can a buffer be considered RO by the compiler after static analisys?

[dipak]

The read-only memory being referred in this context is constant memory so, 1st option is the right one.