I want an array variable to have a program scope.

One way I can do this by passing it as a function pointer throughout the program, which might be complex when we have multiple functions reading/writing this array variable.
Second way to do this, is to have a global variable having program scope. As per the OpenCL specification, Global variables are declared in the program source with the __constant qualifier and are accessed as read-only variables.

//I am writing one sample program to demonstrate my problem:
__constant uint arr[2] = {0, 0}; // an array of unsigned integer
void func1 (uint tmp)
{
for(int i = 0; i < 2; i ++)
arr[i] = tmp+i;
}
void func2(uint tmp)
{
for(int i = 0; i < 2; i++)
tmp = arr[i];
}
__kernel void demoKernel(__global uint *input,
__global uint *output)
{
uint index = get_global_id(0);
func1(input[index]);
func2(output[index]);
}

when i compiled this i got the following error:

tmp/OCLrnlEIO.cl", line 5: error: expression must be a modifiable lvalue
arr[i] = tmp+i;
I searched in google for this error, i found that it is because of type of "arr" is array of 2 length (it is not a pointer).

So my questions are:

1. What is reason for this error and How I can fix this ?
2. secondly, my requirement is to not only to read the array, but also write on it, so how should I use __constant qualifier for that which is read-only variable?

Thanks in Advance !!