Global device variable OpenCL

Question

How can I declare a device variable that is global to all the threads in OpenCL? I'm porting some code from CUDA to OpenCL. In my CUDA implementation I have something like

...
...

__device__ int d_var;

...
...
void foo() {
   int h_var = 0
   cudaMemcpyToSymbol(d_var, h_var, sizeof(int));
   do{
      //launch kernel, inside kernel d_var is modified
      cudaMemcpyFromSymbol(h_var, d_var, sizeof(int));
   }while(h_var != 0);
}

I've been reading through OpenCL example codes but cannot figure out how to do this. Any advise would be great !

Christian Rau · Accepted Answer · 2013-05-14 15:40:49Z

1

Unfortunately this is not that straight-forward to port. Whereas I'm not completely sure if you can at least define and use (read/write) such a global variable in an OpenCL program (but I don't think so), there is definitely no way to access (read/write) this variable from the host.

What you will have to do is put it into the kernel as an additional kernel argument. If it is only writen by the host and read by the kernel an ordinary int variable would suffice (thus a copy for each kernel). But in your case as it is also written in the kernel and read from the host, it has to be a __global pointer, to which an appropriate buffer is bound:

__kernel void func(..., __global int *d_var) { ... }

Which you can then read from and write to on the host using the usual buffer functionality (and in the kernel through classic pointer dereference *d_var, of course, but keep in mind that concurrent writes to this variable have unspecified results if not atomic).

answered May 14, 2013 at 15:40

Christian Rau

46.2k11 gold badges116 silver badges193 bronze badges

Sign up to request clarification or add additional context in comments.

4 Comments

BRabbit27 Over a year ago

Sorry I made a mistake in my code. There is a host h_var and a device d_var and then I copy back and forth from device2host. I think I will have to create another buffer as you suggest. I tried with __kernel void foo(..., __constant int d_var) but I got an error, is it because the d_var is not a pointer?

Christian Rau Over a year ago

@BRabbit27 "is it because the d_var is not a pointer?" - Yes, it is. I understood your code example, the answer still applies. You have to create a buffer for your d_var (being just a single int large, or a sub-buffer of a larger buffer) and copy from there to h_var and vice-versa using the usual buffer read/write operations. Then you pass this buffer as kernel argument for the *d_var parameter, as usual. By the way, if you want to write to d_var in the kernel, as you have written in the question, you need __global rather than __constant.

BRabbit27 Over a year ago

I have to put it as __global inside the Kernel? What if I don't put the qualifier? Its converted into a local variable? Tried both ways but none of them worked. Well, my loop only runs once but I'm sure it has to run more times.

BRabbit27 Over a year ago

My mistake again. I forgot to dereference the pointer inside the kernel. *d_change = 1 I had it d_change = 1. Works perfect :)

Dithermaster · Accepted Answer · 2013-05-15 02:25:39Z

0

For small items that fit in a variable, pass them as arguments.

For larger constant items, use OpenCL "constant" address space. It is optimized for broadcast usage.

In your CUDA example, you read the data back out. In OpenCL you'd need to use a regular buffer object for this, since parameters are not read out back to the host, and constant memory is read-only.

answered May 15, 2013 at 2:25

Dithermaster

6,3531 gold badge15 silver badges20 bronze badges

Collectives™ on Stack Overflow

Global device variable OpenCL

2 Answers 2

4 Comments

Comments

Your Answer

Hot Network Questions

Collectives™ on Stack Overflow

2 Answers 2

4 Comments

Comments

Your Answer

Sign up or log in

Post as a guest

Related