Hello.

Is it possible to have global variables in OpenCL similar to __device__ variables in Cuda which can be assigned via cudaMemcpyToSymbol?

In OpenCL you have 4 address spaces (OpenCL 1.1):

• __private - for every work item
• __local - for all work items inside work group
• __global - for all work items inside whole NDRange
• __constant - for global read-only variables

In addition OpenCL 2.0 (pass in -cl-std=CL2.0 into clBuildProgram), which allows for variable in program scope and static variables to be declared in __global address space. They may be initialized with constant expressions only in code.

Thanks, I have already figured out that the only way in OpenCL to transfer data from host to device is through function parameters. There is no analog of cudaMemcpyToSymbol in OpenCL.

Not really - the only way to write/read/copy buffers between host and device are:

clEnqueueCopyBuffer - one cl_mem buffer to another cl_mem buffer (e.g. copy on device)

clEnqueueReadBuffer - device cl_mem buffer to host memory (e.g. device -> host)

clEnqueueWriteBuffer - host memory to cl_mem buffer (e.g. host -> device)

clEnqueueCopyImageToBuffer - image object into cl_mem buffer object

clEnqueueCopyBufferToImage - cl_mem buffer object to image object

clEnqueueCopyImage - image object to image object

clEnqueueReadImage - image object to host memory

clEnqueueWriteImage - host memory to image object

These are the commands that enqueue read/write/copy operations into command queue. Additionally you can:

clEnqueueMapBuffer - maps region of cl_mem buffer object to host memory

clEnqueueMapImage - ditto for image object

clEnqueueUnmapMemObject - unmaps previously mapped object (either cl_mem or image object)

Compared to CUDA you will however have to pass in pointers as kernel arguments (their memory can be initialized in advance of course). Kernel arguments don't usually hold too much data (mostly just few variables, and pointers to memory that is going to be used).

Out of curiosity (as not many people use OpenCL) - if you can share - what are you trying to do?

In Cuda we can declare a variable this way - __device__ int someVar;

Then via cudaMemcpyToSymbol we can copy/transfer a value to that variable. And we can refer to this variable from kernels. This variable will be available through the whole app life. This is some sort of static variable of TU and this is really handy beasue we don't need to pass such variables as kernel arguments. I see now these is no such thing in OpenCL.

I'm not doing anything specific, just learning OpenCL/Cuda and facinated by their opportunities.

This is the downside of OpenCL, it needs to have most of the data you use passed as kernel arguments. The constant buffers can only be used for pre-defined tables, you don't really write into them.

Most of the time you should be okay with few dozen kernel arguments (I don't think we ever got to some higher numbers like 100). At least not in interactive path tracer (on which I gave up at some point, as it had no commercial product potential at the time) and other projects (which I can't describe).