Global Arrays in OpenCL

Hello,

I am trying to find an example somewhere to let me know if this theory of mine is plausible or not. I have an OpenCL program consisting of multiple kernels that use multiple arrays, several of which are constant throughout the entire run of the program. I am comfortable setting up the arrays and passing them to the device using either


cl_mem clbMemA = clCreateBuffer(oclContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float4)*(size), HostMemA, &oclError);

or

oclError= clEnqueueWriteBuffer(oclCmdQueue, clbMemA, CL_FALSE, 0, sizeof(cl_float4)*(size), HostMemA, 0, NULL, NULL);

I then would use the clSetKernelArg to set the argument for the kernel. This is pretty standard.

However, I am wondering if it is possible to write an OpenCL program that has, generally, the following form:


__global const float2 *Points;

__kernel void CalcSomething(__global float *Values)
{
int i = get_global_id(0);
int j = get_global_id(1);
Values[i] = Points[i].x + Points[j].y;
}

__kernel void CalcSomethingElse(__global float *Values)
{
int i = get_global_id(0);
int j = get_global_id(1);
Values[j] = Points[j].x+Points[i].y;
}

Obviously, the code doesn’t do anything important -it is just a sample, but can I declare “Points” outside of the kernel, and then use either clEnqueueWriteBuffer to store an array (allocated on the host) there? If so, how would I go about doing that?

The reason I ask is that I have four or five kernels that use the same five or six constant arrays throughout the course of obtaining a solution, and it would be convenient to only have to set them up once, instead of as an additional argument for every kernel.

Thank you in advance!

However, I am wondering if it is possible to write an OpenCL program that has, generally, the following form

I’m afraid that’s not allowed by the spec. Variables at global scope have to be constants.

That would explain why I could not find any examples :slight_smile:

Thank you again!